1use runmat_accelerate_api::{GpuTensorHandle, HostTensorView};
4use runmat_builtins::{Tensor, Value};
5use runmat_macros::runtime_builtin;
6
7use super::common::{build_strides, materialize_value, parse_dims};
8use crate::builtins::common::spec::{
9 BroadcastSemantics, BuiltinFusionSpec, BuiltinGpuSpec, ConstantStrategy, GpuOpKind,
10 ProviderHook, ReductionNaN, ResidencyPolicy, ScalarType, ShapeRequirements,
11};
12use crate::builtins::common::tensor;
13#[cfg(feature = "doc_export")]
14use crate::register_builtin_doc_text;
15use crate::{register_builtin_fusion_spec, register_builtin_gpu_spec};
16
17#[cfg(feature = "doc_export")]
18pub const DOC_MD: &str = r#"---
19title: "sub2ind"
20category: "array/indexing"
21keywords: ["sub2ind", "linear index", "column major", "gpu indexing", "nd indexing"]
22summary: "Convert N-D subscripts into MATLAB-style column-major linear indices."
23references: []
24gpu_support:
25 elementwise: false
26 reduction: false
27 precisions: ["f32", "f64"]
28 broadcasting: "matlab"
29 notes: "When a GPU provider exposes the `sub2ind` hook (WGPU today), the conversion runs on the device with bounds and integrality checks; other providers fall back to the host implementation and reupload the result."
30fusion:
31 elementwise: false
32 reduction: false
33 max_inputs: 0
34 constants: "inline"
35requires_feature: null
36tested:
37 unit: "builtins::array::indexing::sub2ind::tests"
38 integration: "builtins::array::indexing::sub2ind::tests::sub2ind_gpu_roundtrip"
39---
40
41# What does the `sub2ind` function do in MATLAB / RunMat?
42`sub2ind(sz, s1, s2, ...)` converts row/column (or higher-dimensional) subscripts into MATLAB's column-major linear indexing form. The size vector `sz` defines the extents of the target array, and you must supply one subscript array per dimension.
43
44## How does the `sub2ind` function behave in MATLAB / RunMat?
45- Subscripts can be scalars or arrays. When arrays are provided, they must share the same size. Scalars broadcast to that common shape.
46- All subscripts must be positive integers within the corresponding dimension's range.
47- The size vector can be a row or column vector. Each element must be a positive integer.
48- Complex, NaN, or infinite values are rejected.
49- The result uses the same shape as the subscript arrays. Scalars produce a scalar double.
50- When any input is a GPU tensor, RunMat computes on the host (to reuse integer semantics) and uploads the resulting indices back to the GPU so fusion and downstream kernels keep operating on device.
51
52## `sub2ind` Function GPU Execution Behaviour
53When a WGPU-backed provider is active, `sub2ind` executes entirely on the GPU. The shader mirrors MATLAB's validation rules: it rejects non-finite values, non-integer subscripts, and out-of-range indices, surfacing the same diagnostic messages as the CPU path. Providers that do not yet implement the hook fall back to the host implementation; after the indices are computed they are uploaded back to the active provider so downstream fused kernels continue operating on device data.
54
55## Examples of using the `sub2ind` function in MATLAB / RunMat
56
57### Converting a single matrix subscript to a linear index
58
59```matlab
60idx = sub2ind([3 4], 2, 3);
61```
62
63Expected output:
64
65```matlab
66idx = 8;
67```
68
69### Mapping multiple subscripts into one-dimensional indices
70
71```matlab
72rows = [1; 2; 3];
73cols = [3; 3; 3];
74idx = sub2ind([3 5], rows, cols);
75```
76
77Expected output:
78
79```matlab
80idx =
81 7
82 8
83 9
84```
85
86### Handling higher-dimensional array subscripts
87
88```matlab
89row = [1 1];
90col = [2 3];
91page = [1 2];
92idx = sub2ind([2 3 4], row, col, page);
93```
94
95Expected output:
96
97```matlab
98idx = [3 11];
99```
100
101### Broadcasting scalar subscripts across array inputs
102
103```matlab
104rows = [1 2 3];
105idx = sub2ind([3 4], rows, 4);
106```
107
108Expected output:
109
110```matlab
111idx = [10 11 12];
112```
113
114### Retaining GPU residency for batched index conversions
115
116```matlab
117rows = gpuArray((1:100)');
118cols = gpuArray(ones(100, 1) * 4);
119idx = sub2ind([100 4], rows, cols);
120```
121
122Expected behavior:
123
124```matlab
125% idx remains a gpuArray containing the column-major indices.
126disp(gather(idx(1:5)));
127% Output:
128% 301
129% 302
130% 303
131% 304
132% 305
133```
134
135### Detecting invalid out-of-range subscripts
136
137```matlab
138try
139 idx = sub2ind([3 4], 4, 1);
140catch ME
141 disp(ME.message);
142end
143```
144
145Expected output:
146
147```matlab
148Index exceeds the number of rows in dimension 1.
149```
150
151## GPU residency in RunMat (Do I need `gpuArray`?)
152You typically do **not** need to call `gpuArray` yourself. When the active provider implements the `sub2ind` hook (WGPU today), the entire conversion runs on the GPU and returns a device tensor. If no provider is available, or the provider lacks the hook, RunMat falls back to the host implementation and uploads the resulting indices back to the GPU so residency is maintained automatically.
153
154## FAQ
155
156### What data types does `sub2ind` accept?
157Numeric and logical inputs are accepted. Logical values are converted to doubles before validation. Complex, NaN, and infinite values are rejected with a descriptive error.
158
159### Can the size vector contain zeros?
160No. Every dimension size must be a positive integer. This matches MATLAB's behavior for index conversion.
161
162### Do subscripts have to be the same size?
163Yes. All non-scalar subscripts must share the same size (shape). Scalars broadcast to that common shape.
164
165### What happens when subscripts are out of range?
166`sub2ind` throws an error explaining which dimension failed the bounds check. This mirrors MATLAB's run-time error.
167
168### Does the function support GPU arrays?
169Yes. With the WGPU provider the conversion happens entirely on device, including validation. Other providers gather the data to the host, compute the indices, and upload them back to the device automatically.
170
171### Are fractional subscripts rounded?
172No. Non-integer, NaN, or infinite subscripts raise an error.
173
174### How is the linear index computed?
175The output uses MATLAB's column-major convention: `1 + sum((s_k - 1) * stride_k)` where `stride_k` is the product of the preceding dimensions.
176
177### Can I call `sub2ind` with more subscripts than dimensions?
178No. You must pass exactly one subscript per dimension listed in the size vector.
179
180### What about empty outputs?
181If the subscript arrays are empty, `sub2ind` returns an empty double array with the same shape.
182
183### Does `sub2ind` change the orientation of row/column vectors?
184No. The output preserves the orientation (shape) of the subscript arrays, so row vectors stay row vectors and column vectors stay column vectors.
185
186## See Also
187[ind2sub](./ind2sub), [find](./find), [size](../../introspection/size), [gpuArray](../../acceleration/gpu/gpuArray), [gather](../../acceleration/gpu/gather)
188
189## Source & Feedback
190- The full source code for the implementation of the `sub2ind` function is available at: [`crates/runmat-runtime/src/builtins/array/indexing/sub2ind.rs`](https://github.com/runmat-org/runmat/blob/main/crates/runmat-runtime/src/builtins/array/indexing/sub2ind.rs)
191- Found a bug or behavioral difference? Please [open an issue](https://github.com/runmat-org/runmat/issues/new/choose) with details and a minimal repro.
192"#;
193
194pub const GPU_SPEC: BuiltinGpuSpec = BuiltinGpuSpec {
195 name: "sub2ind",
196 op_kind: GpuOpKind::Custom("indexing"),
197 supported_precisions: &[ScalarType::F32, ScalarType::F64],
198 broadcast: BroadcastSemantics::Matlab,
199 provider_hooks: &[ProviderHook::Custom("sub2ind")],
200 constant_strategy: ConstantStrategy::InlineLiteral,
201 residency: ResidencyPolicy::NewHandle,
202 nan_mode: ReductionNaN::Include,
203 two_pass_threshold: None,
204 workgroup_size: None,
205 accepts_nan_mode: false,
206 notes: "Providers can implement the custom `sub2ind` hook to execute on device; runtimes fall back to host computation otherwise.",
207};
208
209register_builtin_gpu_spec!(GPU_SPEC);
210
211pub const FUSION_SPEC: BuiltinFusionSpec = BuiltinFusionSpec {
212 name: "sub2ind",
213 shape: ShapeRequirements::Any,
214 constant_strategy: ConstantStrategy::InlineLiteral,
215 elementwise: None,
216 reduction: None,
217 emits_nan: false,
218 notes: "Index conversion executes eagerly on the host; fusion does not apply.",
219};
220
221register_builtin_fusion_spec!(FUSION_SPEC);
222
223#[cfg(feature = "doc_export")]
224register_builtin_doc_text!("sub2ind", DOC_MD);
225
226#[runtime_builtin(
227 name = "sub2ind",
228 category = "array/indexing",
229 summary = "Convert N-D subscripts into MATLAB-style column-major linear indices.",
230 keywords = "sub2ind,linear index,column major,gpu indexing",
231 accel = "custom"
232)]
233fn sub2ind_builtin(dims_val: Value, rest: Vec<Value>) -> Result<Value, String> {
234 let (dims_value, dims_was_gpu) = materialize_value(dims_val)?;
235 let dims = parse_dims(&dims_value)?;
236 if dims.is_empty() {
237 return Err("Size vector must have at least one element.".to_string());
238 }
239
240 if rest.len() != dims.len() {
241 return Err("The number of subscripts supplied must equal the number of dimensions in the size vector.".to_string());
242 }
243
244 if let Some(value) = try_gpu_sub2ind(&dims, &rest)? {
245 return Ok(value);
246 }
247
248 let mut saw_gpu = dims_was_gpu;
249 let mut subscripts: Vec<Tensor> = Vec::with_capacity(rest.len());
250 for value in rest {
251 let (materialised, was_gpu) = materialize_value(value)?;
252 saw_gpu |= was_gpu;
253 let tensor = tensor::value_into_tensor_for("sub2ind", materialised)?;
254 subscripts.push(tensor);
255 }
256
257 let (result_data, result_shape) = compute_indices(&dims, &subscripts)?;
258 let want_gpu_output = saw_gpu && runmat_accelerate_api::provider().is_some();
259
260 if want_gpu_output {
261 #[cfg(all(test, feature = "wgpu"))]
262 {
263 if runmat_accelerate_api::provider().is_none() {
264 let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
265 runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
266 );
267 }
268 }
269 let shape = result_shape.clone().unwrap_or_else(|| vec![1, 1]);
270 if let Some(provider) = runmat_accelerate_api::provider() {
271 let view = HostTensorView {
272 data: &result_data,
273 shape: &shape,
274 };
275 if let Ok(handle) = provider.upload(&view) {
276 return Ok(Value::GpuTensor(handle));
277 }
278 }
279 }
280
281 build_host_value(result_data, result_shape)
282}
283
284fn try_gpu_sub2ind(dims: &[usize], subs: &[Value]) -> Result<Option<Value>, String> {
285 #[cfg(all(test, feature = "wgpu"))]
286 {
287 if subs
288 .iter()
289 .any(|v| matches!(v, Value::GpuTensor(h) if h.device_id != 0))
290 {
291 let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
292 runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
293 );
294 }
295 }
296 let provider = match runmat_accelerate_api::provider() {
297 Some(p) => p,
298 None => return Ok(None),
299 };
300 if !subs
301 .iter()
302 .all(|value| matches!(value, Value::GpuTensor(_)))
303 {
304 return Ok(None);
305 }
306 if dims.is_empty() {
307 return Ok(None);
308 }
309
310 let mut handles: Vec<&GpuTensorHandle> = Vec::with_capacity(subs.len());
311 for value in subs {
312 if let Value::GpuTensor(handle) = value {
313 handles.push(handle);
314 }
315 }
316
317 if handles.len() != dims.len() {
318 return Err("The number of subscripts supplied must equal the number of dimensions in the size vector.".to_string());
319 }
320
321 let mut scalar_mask: Vec<bool> = Vec::with_capacity(handles.len());
322 let mut target_shape: Option<Vec<usize>> = None;
323 let mut result_len: usize = 1;
324 let mut saw_non_scalar = false;
325
326 for handle in &handles {
327 let len = tensor::element_count(&handle.shape);
328 let is_scalar = len == 1;
329 scalar_mask.push(is_scalar);
330 if !is_scalar {
331 saw_non_scalar = true;
332 if let Some(existing) = &target_shape {
333 if existing != &handle.shape {
334 return Err("Subscript inputs must have the same size.".to_string());
335 }
336 } else {
337 target_shape = Some(handle.shape.clone());
338 result_len = len;
339 }
340 }
341 }
342
343 if !saw_non_scalar {
344 target_shape = Some(vec![1, 1]);
345 result_len = 1;
346 } else if let Some(shape) = &target_shape {
347 result_len = tensor::element_count(shape);
348 }
349
350 let strides = build_strides(dims)?;
351 if dims.iter().any(|&d| d > u32::MAX as usize)
352 || strides.iter().any(|&s| s > u32::MAX as usize)
353 || result_len > u32::MAX as usize
354 {
355 return Ok(None);
356 }
357
358 let output_shape = target_shape.clone().unwrap_or_else(|| vec![1, 1]);
359 match provider.sub2ind(
360 dims,
361 &strides,
362 &handles,
363 &scalar_mask,
364 result_len,
365 &output_shape,
366 ) {
367 Ok(handle) => Ok(Some(Value::GpuTensor(handle))),
368 Err(err) => Err(err.to_string()),
369 }
370}
371
372fn compute_indices(
373 dims: &[usize],
374 subscripts: &[Tensor],
375) -> Result<(Vec<f64>, Option<Vec<usize>>), String> {
376 let mut target_shape: Option<Vec<usize>> = None;
377 let mut result_len: usize = 1;
378 let mut has_non_scalar = false;
379
380 for tensor in subscripts {
381 if tensor.data.len() != 1 {
382 has_non_scalar = true;
383 if let Some(shape) = &target_shape {
384 if &tensor.shape != shape {
385 return Err("Subscript inputs must have the same size.".to_string());
386 }
387 } else {
388 target_shape = Some(tensor.shape.clone());
389 result_len = tensor.data.len();
390 }
391 }
392 }
393
394 if !has_non_scalar {
395 target_shape = Some(vec![1, 1]);
397 result_len = 1;
398 }
399
400 if result_len == 0 {
401 return Ok((Vec::new(), target_shape));
402 }
403
404 let strides = build_strides(dims)?;
405 let mut output = Vec::with_capacity(result_len);
406
407 for idx in 0..result_len {
408 let mut offset: usize = 0;
409 for (dim_index, (&dim, tensor)) in dims.iter().zip(subscripts.iter()).enumerate() {
410 let raw = subscript_value(tensor, idx);
411 let coerced = coerce_subscript(raw, dim_index + 1, dim)?;
412 let term = coerced
413 .checked_sub(1)
414 .and_then(|v| v.checked_mul(strides[dim_index]))
415 .ok_or_else(|| "Index exceeds array dimensions.".to_string())?;
416 offset = offset
417 .checked_add(term)
418 .ok_or_else(|| "Index exceeds array dimensions.".to_string())?;
419 }
420 output.push((offset + 1) as f64);
421 }
422
423 Ok((output, target_shape))
424}
425
426fn subscript_value(tensor: &Tensor, idx: usize) -> f64 {
427 if tensor.data.len() == 1 {
428 tensor.data[0]
429 } else {
430 tensor.data[idx]
431 }
432}
433
434fn coerce_subscript(value: f64, dim_number: usize, dim_size: usize) -> Result<usize, String> {
435 if !value.is_finite() {
436 return Err(
437 "Subscript indices must either be real positive integers or logicals.".to_string(),
438 );
439 }
440 let rounded = value.round();
441 if (rounded - value).abs() > f64::EPSILON {
442 return Err(
443 "Subscript indices must either be real positive integers or logicals.".to_string(),
444 );
445 }
446 if rounded < 1.0 {
447 return Err(
448 "Subscript indices must either be real positive integers or logicals.".to_string(),
449 );
450 }
451 if rounded > dim_size as f64 {
452 return Err(dimension_bounds_error(dim_number));
453 }
454 Ok(rounded as usize)
455}
456
457fn dimension_bounds_error(dim_number: usize) -> String {
458 match dim_number {
459 1 => format!("Index exceeds the number of rows in dimension {dim_number}."),
460 2 => format!("Index exceeds the number of columns in dimension {dim_number}."),
461 3 => format!("Index exceeds the number of pages in dimension {dim_number}."),
462 _ => "Index exceeds array dimensions.".to_string(),
463 }
464}
465
466fn build_host_value(data: Vec<f64>, shape: Option<Vec<usize>>) -> Result<Value, String> {
467 let shape = shape.unwrap_or_else(|| vec![1, 1]);
468 if data.len() == 1 && tensor::element_count(&shape) == 1 {
469 Ok(Value::Num(data[0]))
470 } else {
471 let tensor = Tensor::new(data, shape)
472 .map_err(|e| format!("Unable to construct sub2ind output: {e}"))?;
473 Ok(Value::Tensor(tensor))
474 }
475}
476
477#[cfg(test)]
478mod tests {
479 use super::*;
480 use crate::builtins::common::test_support;
481 use runmat_builtins::{IntValue, Tensor, Value};
482
483 #[test]
484 fn converts_scalar_indices() {
485 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
486 let result =
487 sub2ind_builtin(Value::Tensor(dims), vec![Value::Num(2.0), Value::Num(3.0)]).unwrap();
488 assert_eq!(result, Value::Num(8.0));
489 }
490
491 #[test]
492 fn broadcasts_scalars_over_vectors() {
493 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
494 let rows = Tensor::new(vec![1.0, 2.0, 3.0], vec![3, 1]).unwrap();
495 let result = sub2ind_builtin(
496 Value::Tensor(dims),
497 vec![Value::Tensor(rows), Value::Num(4.0)],
498 )
499 .unwrap();
500 match result {
501 Value::Tensor(t) => {
502 assert_eq!(t.shape, vec![3, 1]);
503 assert_eq!(t.data, vec![10.0, 11.0, 12.0]);
504 }
505 other => panic!("expected tensor result, got {other:?}"),
506 }
507 }
508
509 #[test]
510 fn handles_three_dimensions() {
511 let dims = Tensor::new(vec![2.0, 3.0, 4.0], vec![1, 3]).unwrap();
512 let row = Tensor::new(vec![1.0, 1.0], vec![1, 2]).unwrap();
513 let col = Tensor::new(vec![2.0, 3.0], vec![1, 2]).unwrap();
514 let page = Tensor::new(vec![1.0, 2.0], vec![1, 2]).unwrap();
515 let result = sub2ind_builtin(
516 Value::Tensor(dims),
517 vec![Value::Tensor(row), Value::Tensor(col), Value::Tensor(page)],
518 )
519 .unwrap();
520 match result {
521 Value::Tensor(t) => {
522 assert_eq!(t.shape, vec![1, 2]);
523 assert_eq!(t.data, vec![3.0, 11.0]);
524 }
525 other => panic!("expected tensor result, got {other:?}"),
526 }
527 }
528
529 #[test]
530 fn rejects_out_of_range_subscripts() {
531 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
532 let err = sub2ind_builtin(Value::Tensor(dims), vec![Value::Num(4.0), Value::Num(1.0)])
533 .unwrap_err();
534 assert!(
535 err.contains("Index exceeds"),
536 "expected index bounds error, got {err}"
537 );
538 }
539
540 #[test]
541 fn rejects_shape_mismatch() {
542 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
543 let rows = Tensor::new(vec![1.0, 2.0], vec![2, 1]).unwrap();
544 let cols = Tensor::new(vec![1.0, 2.0, 3.0], vec![3, 1]).unwrap();
545 let err = sub2ind_builtin(
546 Value::Tensor(dims),
547 vec![Value::Tensor(rows), Value::Tensor(cols)],
548 )
549 .unwrap_err();
550 assert!(
551 err.contains("same size"),
552 "expected size mismatch error, got {err}"
553 );
554 }
555
556 #[test]
557 fn rejects_non_integer_subscripts() {
558 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
559 let err = sub2ind_builtin(Value::Tensor(dims), vec![Value::Num(1.5), Value::Num(1.0)])
560 .unwrap_err();
561 assert!(
562 err.contains("real positive integers"),
563 "expected integer coercion error, got {err}"
564 );
565 }
566
567 #[test]
568 fn accepts_integer_value_variants() {
569 let dims = Value::Tensor(Tensor::new(vec![3.0], vec![1, 1]).unwrap());
570 let result = sub2ind_builtin(dims, vec![Value::Int(IntValue::I32(2))]).expect("sub2ind");
571 assert_eq!(result, Value::Num(2.0));
572 }
573
574 #[test]
575 fn sub2ind_gpu_roundtrip() {
576 test_support::with_test_provider(|provider| {
577 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
578 let rows = Tensor::new(vec![1.0, 2.0, 3.0], vec![3, 1]).unwrap();
579 let cols = Tensor::new(vec![4.0, 4.0, 4.0], vec![3, 1]).unwrap();
580
581 let dims_handle = provider
582 .upload(&HostTensorView {
583 data: &dims.data,
584 shape: &dims.shape,
585 })
586 .expect("upload dims");
587 let rows_handle = provider
588 .upload(&HostTensorView {
589 data: &rows.data,
590 shape: &rows.shape,
591 })
592 .expect("upload rows");
593 let cols_handle = provider
594 .upload(&HostTensorView {
595 data: &cols.data,
596 shape: &cols.shape,
597 })
598 .expect("upload cols");
599
600 let result = sub2ind_builtin(
601 Value::GpuTensor(dims_handle),
602 vec![Value::GpuTensor(rows_handle), Value::GpuTensor(cols_handle)],
603 )
604 .expect("sub2ind");
605
606 match result {
607 Value::GpuTensor(handle) => {
608 let gathered = test_support::gather(Value::GpuTensor(handle)).unwrap();
609 assert_eq!(gathered.shape, vec![3, 1]);
610 assert_eq!(gathered.data, vec![10.0, 11.0, 12.0]);
611 }
612 other => panic!("expected gpu tensor, got {other:?}"),
613 }
614 });
615 }
616
617 #[test]
618 #[cfg(feature = "wgpu")]
619 fn sub2ind_wgpu_matches_cpu() {
620 let _ = runmat_accelerate::backend::wgpu::provider::register_wgpu_provider(
621 runmat_accelerate::backend::wgpu::provider::WgpuProviderOptions::default(),
622 );
623 let Some(provider) = runmat_accelerate_api::provider() else {
624 panic!("wgpu provider not available");
625 };
626
627 let dims = Tensor::new(vec![3.0, 4.0], vec![1, 2]).unwrap();
628 let rows = Tensor::new(vec![1.0, 2.0, 3.0], vec![3, 1]).unwrap();
629 let cols = Tensor::new(vec![4.0, 4.0, 4.0], vec![3, 1]).unwrap();
630
631 let cpu = sub2ind_builtin(
632 Value::Tensor(dims.clone()),
633 vec![Value::Tensor(rows.clone()), Value::Tensor(cols.clone())],
634 )
635 .expect("cpu sub2ind");
636
637 let rows_handle = provider
638 .upload(&HostTensorView {
639 data: &rows.data,
640 shape: &rows.shape,
641 })
642 .expect("upload rows");
643 let cols_handle = provider
644 .upload(&HostTensorView {
645 data: &cols.data,
646 shape: &cols.shape,
647 })
648 .expect("upload cols");
649
650 let result = sub2ind_builtin(
651 Value::Tensor(dims),
652 vec![Value::GpuTensor(rows_handle), Value::GpuTensor(cols_handle)],
653 )
654 .expect("wgpu sub2ind");
655
656 let gathered = test_support::gather(result).expect("gather");
657 let expected = match cpu {
658 Value::Tensor(t) => t,
659 Value::Num(v) => Tensor::new(vec![v], vec![1, 1]).unwrap(),
660 other => panic!("unexpected cpu result {other:?}"),
661 };
662 assert_eq!(gathered.shape, expected.shape);
663 assert_eq!(gathered.data, expected.data);
664 }
665
666 #[test]
667 #[cfg(feature = "doc_export")]
668 fn doc_examples_present() {
669 let blocks = test_support::doc_examples(DOC_MD);
670 assert!(!blocks.is_empty());
671 }
672}