1#[cfg(feature = "python")]
2use numpy::{IntoPyArray, PyArray1, PyArrayMethods, PyReadonlyArray1};
3#[cfg(feature = "python")]
4use pyo3::exceptions::PyValueError;
5#[cfg(feature = "python")]
6use pyo3::prelude::*;
7
8#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
9use wasm_bindgen::prelude::*;
10
11use crate::utilities::data_loader::{source_type, Candles};
12use crate::utilities::enums::Kernel;
13use crate::utilities::helpers::{
14 alloc_with_nan_prefix, detect_best_batch_kernel, detect_best_kernel, init_matrix_prefixes,
15 make_uninit_matrix,
16};
17#[cfg(feature = "python")]
18use crate::utilities::kernel_validation::validate_kernel;
19#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
20use core::arch::x86_64::*;
21use thiserror::Error;
22
23#[derive(Debug, Clone)]
24pub enum NviData<'a> {
25 Candles {
26 candles: &'a Candles,
27 close_source: &'a str,
28 },
29 Slices {
30 close: &'a [f64],
31 volume: &'a [f64],
32 },
33}
34
35#[derive(Debug, Clone)]
36pub struct NviOutput {
37 pub values: Vec<f64>,
38}
39
40#[derive(Debug, Clone, Default)]
41pub struct NviParams;
42
43#[derive(Debug, Clone)]
44pub struct NviInput<'a> {
45 pub data: NviData<'a>,
46 pub params: NviParams,
47}
48
49impl<'a> NviInput<'a> {
50 #[inline]
51 pub fn from_candles(candles: &'a Candles, close_source: &'a str, params: NviParams) -> Self {
52 Self {
53 data: NviData::Candles {
54 candles,
55 close_source,
56 },
57 params,
58 }
59 }
60 #[inline]
61 pub fn from_slices(close: &'a [f64], volume: &'a [f64], params: NviParams) -> Self {
62 Self {
63 data: NviData::Slices { close, volume },
64 params,
65 }
66 }
67 #[inline]
68 pub fn with_default_candles(candles: &'a Candles) -> Self {
69 Self::from_candles(candles, "close", NviParams)
70 }
71}
72
73#[derive(Debug, Error)]
74pub enum NviError {
75 #[error("nvi: Empty data provided.")]
76 EmptyInputData,
77 #[error("nvi: Empty data provided.")]
78 EmptyData,
79 #[error("nvi: All values are NaN in both close and volume.")]
80 AllValuesNaN,
81 #[error("nvi: All close values are NaN.")]
82 AllCloseValuesNaN,
83 #[error("nvi: All volume values are NaN.")]
84 AllVolumeValuesNaN,
85 #[error("nvi: Not enough valid data: needed = {needed}, valid = {valid}")]
86 NotEnoughValidData { needed: usize, valid: usize },
87 #[error("nvi: Close and volume length mismatch: close={close_len}, volume={volume_len}")]
88 MismatchedLength { close_len: usize, volume_len: usize },
89 #[error(
90 "nvi: Destination length mismatch: dst={dst_len}, close={close_len}, volume={volume_len}"
91 )]
92 DestinationLengthMismatch {
93 dst_len: usize,
94 close_len: usize,
95 volume_len: usize,
96 },
97 #[error("nvi: output length mismatch: expected = {expected}, got = {got}")]
98 OutputLengthMismatch { expected: usize, got: usize },
99 #[error("nvi: invalid range: start={start}, end={end}, step={step}")]
100 InvalidRange {
101 start: usize,
102 end: usize,
103 step: usize,
104 },
105 #[error("nvi: invalid kernel for batch: {0:?}")]
106 InvalidKernelForBatch(Kernel),
107 #[error("nvi: invalid period: period = {period}, data length = {data_len}")]
108 InvalidPeriod { period: usize, data_len: usize },
109}
110
111#[derive(Copy, Clone, Debug, Default)]
112pub struct NviBuilder {
113 kernel: Kernel,
114}
115impl NviBuilder {
116 #[inline(always)]
117 pub fn new() -> Self {
118 Self {
119 kernel: Kernel::Auto,
120 }
121 }
122 #[inline(always)]
123 pub fn kernel(mut self, k: Kernel) -> Self {
124 self.kernel = k;
125 self
126 }
127 #[inline(always)]
128 pub fn apply(self, c: &Candles) -> Result<NviOutput, NviError> {
129 let i = NviInput::with_default_candles(c);
130 nvi_with_kernel(&i, self.kernel)
131 }
132 #[inline(always)]
133 pub fn apply_slice(self, close: &[f64], volume: &[f64]) -> Result<NviOutput, NviError> {
134 let i = NviInput::from_slices(close, volume, NviParams);
135 nvi_with_kernel(&i, self.kernel)
136 }
137 #[inline(always)]
138 pub fn into_stream(self) -> Result<NviStream, NviError> {
139 NviStream::try_new()
140 }
141}
142
143#[derive(Debug, Clone)]
144pub struct NviStream {
145 prev_close: f64,
146 prev_volume: f64,
147 nvi_val: f64,
148 started: bool,
149}
150
151impl NviStream {
152 #[inline]
153 pub fn try_new() -> Result<Self, NviError> {
154 Ok(Self {
155 prev_close: 0.0,
156 prev_volume: 0.0,
157 nvi_val: 1000.0,
158 started: false,
159 })
160 }
161
162 #[inline(always)]
163 pub fn update(&mut self, close: f64, volume: f64) -> Option<f64> {
164 if !self.started {
165 if close.is_nan() || volume.is_nan() {
166 return None;
167 }
168 self.prev_close = close;
169 self.prev_volume = volume;
170 self.started = true;
171 return Some(self.nvi_val);
172 }
173
174 let mut nvi = self.nvi_val;
175 if volume < self.prev_volume {
176 let pct = (close - self.prev_close) / self.prev_close;
177 nvi += nvi * pct;
178 }
179
180 self.nvi_val = nvi;
181 self.prev_close = close;
182 self.prev_volume = volume;
183
184 Some(nvi)
185 }
186}
187
188#[derive(Clone, Debug)]
189pub struct NviBatchOutput {
190 pub values: Vec<f64>,
191 pub rows: usize,
192 pub cols: usize,
193}
194
195#[inline]
196pub fn nvi(input: &NviInput) -> Result<NviOutput, NviError> {
197 nvi_with_kernel(input, Kernel::Auto)
198}
199pub fn nvi_with_kernel(input: &NviInput, kernel: Kernel) -> Result<NviOutput, NviError> {
200 let (close, volume): (&[f64], &[f64]) = match &input.data {
201 NviData::Candles {
202 candles,
203 close_source,
204 } => {
205 let close = source_type(candles, close_source);
206 let volume = candles
207 .select_candle_field("volume")
208 .map_err(|_| NviError::EmptyInputData)?;
209 (close, volume)
210 }
211 NviData::Slices { close, volume } => (*close, *volume),
212 };
213
214 if close.is_empty() || volume.is_empty() {
215 return Err(NviError::EmptyInputData);
216 }
217 if close.len() != volume.len() {
218 return Err(NviError::MismatchedLength {
219 close_len: close.len(),
220 volume_len: volume.len(),
221 });
222 }
223 let first = close
224 .iter()
225 .zip(volume)
226 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
227 .ok_or_else(|| {
228 if close.iter().all(|&c| c.is_nan()) {
229 NviError::AllCloseValuesNaN
230 } else {
231 NviError::AllVolumeValuesNaN
232 }
233 })?;
234 if close.len() - first < 2 {
235 return Err(NviError::NotEnoughValidData {
236 needed: 2,
237 valid: close.len() - first,
238 });
239 }
240 let mut out = alloc_with_nan_prefix(close.len(), first);
241 let mut chosen = match kernel {
242 Kernel::Auto => Kernel::Scalar,
243 other => other,
244 };
245
246 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
247 if matches!(kernel, Kernel::Auto) && matches!(chosen, Kernel::Avx512 | Kernel::Avx512Batch) {
248 chosen = Kernel::Avx2;
249 }
250 unsafe {
251 match chosen {
252 Kernel::Scalar | Kernel::ScalarBatch => nvi_scalar(close, volume, first, &mut out),
253 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
254 Kernel::Avx2 | Kernel::Avx2Batch => nvi_avx2(close, volume, first, &mut out),
255 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
256 Kernel::Avx512 | Kernel::Avx512Batch => nvi_avx512(close, volume, first, &mut out),
257 _ => unreachable!(),
258 }
259 }
260 Ok(NviOutput { values: out })
261}
262
263#[cfg(not(all(target_arch = "wasm32", feature = "wasm")))]
264#[inline]
265pub fn nvi_into(input: &NviInput, out: &mut [f64]) -> Result<(), NviError> {
266 let (close, volume): (&[f64], &[f64]) = match &input.data {
267 NviData::Candles {
268 candles,
269 close_source,
270 } => {
271 let close = source_type(candles, close_source);
272 let volume = candles
273 .select_candle_field("volume")
274 .map_err(|_| NviError::EmptyInputData)?;
275 (close, volume)
276 }
277 NviData::Slices { close, volume } => (*close, *volume),
278 };
279
280 nvi_into_slice(out, close, volume, Kernel::Auto)
281}
282
283#[inline]
284pub fn nvi_into_slice(
285 dst: &mut [f64],
286 close: &[f64],
287 volume: &[f64],
288 kern: Kernel,
289) -> Result<(), NviError> {
290 if close.is_empty() || volume.is_empty() {
291 return Err(NviError::EmptyInputData);
292 }
293 if close.len() != volume.len() {
294 return Err(NviError::MismatchedLength {
295 close_len: close.len(),
296 volume_len: volume.len(),
297 });
298 }
299 if dst.len() != close.len() {
300 return Err(NviError::OutputLengthMismatch {
301 expected: close.len(),
302 got: dst.len(),
303 });
304 }
305
306 let first = close
307 .iter()
308 .zip(volume)
309 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
310 .ok_or_else(|| {
311 if close.iter().all(|&c| c.is_nan()) {
312 NviError::AllCloseValuesNaN
313 } else {
314 NviError::AllVolumeValuesNaN
315 }
316 })?;
317
318 if close.len() - first < 2 {
319 return Err(NviError::NotEnoughValidData {
320 needed: 2,
321 valid: close.len() - first,
322 });
323 }
324
325 let mut chosen = match kern {
326 Kernel::Auto => Kernel::Scalar,
327 other => other,
328 };
329 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
330 if matches!(kern, Kernel::Auto) && matches!(chosen, Kernel::Avx512 | Kernel::Avx512Batch) {
331 chosen = Kernel::Avx2;
332 }
333
334 unsafe {
335 match chosen {
336 Kernel::Scalar | Kernel::ScalarBatch => nvi_scalar(close, volume, first, dst),
337 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
338 Kernel::Avx2 | Kernel::Avx2Batch => nvi_avx2(close, volume, first, dst),
339 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
340 Kernel::Avx512 | Kernel::Avx512Batch => nvi_avx512(close, volume, first, dst),
341 _ => unreachable!(),
342 }
343 }
344
345 for v in &mut dst[..first] {
346 *v = f64::NAN;
347 }
348
349 Ok(())
350}
351
352pub fn nvi_scalar(close: &[f64], volume: &[f64], first_valid: usize, out: &mut [f64]) {
353 debug_assert!(
354 close.len() == volume.len() && volume.len() == out.len(),
355 "Input slices must all have the same length."
356 );
357
358 let len = close.len();
359 if len == 0 || first_valid >= len {
360 return;
361 }
362
363 let mut nvi_val = 1000.0;
364
365 unsafe {
366 let close_ptr = close.as_ptr();
367 let vol_ptr = volume.as_ptr();
368 let out_ptr = out.as_mut_ptr();
369
370 *out_ptr.add(first_valid) = nvi_val;
371
372 let mut i = first_valid + 1;
373 if i >= len {
374 return;
375 }
376
377 let mut prev_close = *close_ptr.add(i - 1);
378 let mut prev_volume = *vol_ptr.add(i - 1);
379
380 while i < len {
381 let c = *close_ptr.add(i);
382 let v = *vol_ptr.add(i);
383
384 if v < prev_volume {
385 let pct = (c - prev_close) / prev_close;
386 nvi_val += nvi_val * pct;
387 }
388
389 *out_ptr.add(i) = nvi_val;
390
391 prev_close = c;
392 prev_volume = v;
393 i += 1;
394 }
395 }
396}
397
398#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
399#[inline]
400pub unsafe fn nvi_avx2(close: &[f64], volume: &[f64], first_valid: usize, out: &mut [f64]) {
401 let len = close.len();
402 if len == 0 || first_valid >= len {
403 return;
404 }
405
406 let close_ptr = close.as_ptr();
407 let vol_ptr = volume.as_ptr();
408 let out_ptr = out.as_mut_ptr();
409
410 let mut nvi_val = 1000.0;
411 *out_ptr.add(first_valid) = nvi_val;
412
413 let mut i = first_valid + 1;
414 if i >= len {
415 return;
416 }
417
418 while i + 3 < len {
419 let curr_c = _mm256_loadu_pd(close_ptr.add(i) as *const f64);
420 let prev_c = _mm256_loadu_pd(close_ptr.add(i - 1) as *const f64);
421
422 let curr_v = _mm256_loadu_pd(vol_ptr.add(i) as *const f64);
423 let prev_v = _mm256_loadu_pd(vol_ptr.add(i - 1) as *const f64);
424
425 let delta = _mm256_sub_pd(curr_c, prev_c);
426 let pct_raw = _mm256_div_pd(delta, prev_c);
427
428 let mask = _mm256_cmp_pd(curr_v, prev_v, _CMP_LT_OQ);
429 let pct_masked = _mm256_and_pd(pct_raw, mask);
430
431 let mut pcts: [f64; 4] = [0.0; 4];
432 _mm256_storeu_pd(pcts.as_mut_ptr(), pct_masked);
433
434 nvi_val += nvi_val * pcts[0];
435 *out_ptr.add(i) = nvi_val;
436
437 nvi_val += nvi_val * pcts[1];
438 *out_ptr.add(i + 1) = nvi_val;
439
440 nvi_val += nvi_val * pcts[2];
441 *out_ptr.add(i + 2) = nvi_val;
442
443 nvi_val += nvi_val * pcts[3];
444 *out_ptr.add(i + 3) = nvi_val;
445
446 i += 4;
447 }
448
449 while i < len {
450 let c = *close_ptr.add(i);
451 let v = *vol_ptr.add(i);
452
453 if v < *vol_ptr.add(i - 1) {
454 let pct = (c - *close_ptr.add(i - 1)) / *close_ptr.add(i - 1);
455 nvi_val += nvi_val * pct;
456 }
457 *out_ptr.add(i) = nvi_val;
458 i += 1;
459 }
460}
461
462#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
463#[inline]
464pub unsafe fn nvi_avx512(close: &[f64], volume: &[f64], first_valid: usize, out: &mut [f64]) {
465 let len = close.len();
466 if len == 0 || first_valid >= len {
467 return;
468 }
469
470 let close_ptr = close.as_ptr();
471 let vol_ptr = volume.as_ptr();
472 let out_ptr = out.as_mut_ptr();
473
474 let mut nvi_val = 1000.0;
475 *out_ptr.add(first_valid) = nvi_val;
476
477 let mut i = first_valid + 1;
478 if i >= len {
479 return;
480 }
481
482 while i + 7 < len {
483 let curr_c = _mm512_loadu_pd(close_ptr.add(i) as *const f64);
484 let prev_c = _mm512_loadu_pd(close_ptr.add(i - 1) as *const f64);
485
486 let curr_v = _mm512_loadu_pd(vol_ptr.add(i) as *const f64);
487 let prev_v = _mm512_loadu_pd(vol_ptr.add(i - 1) as *const f64);
488
489 let delta = _mm512_sub_pd(curr_c, prev_c);
490 let pct_raw = _mm512_div_pd(delta, prev_c);
491
492 let m = _mm512_cmp_pd_mask(curr_v, prev_v, _CMP_LT_OQ);
493 let pct_masked = _mm512_maskz_mov_pd(m, pct_raw);
494
495 let mut pcts: [f64; 8] = [0.0; 8];
496 _mm512_storeu_pd(pcts.as_mut_ptr(), pct_masked);
497
498 nvi_val += nvi_val * pcts[0];
499 *out_ptr.add(i) = nvi_val;
500 nvi_val += nvi_val * pcts[1];
501 *out_ptr.add(i + 1) = nvi_val;
502 nvi_val += nvi_val * pcts[2];
503 *out_ptr.add(i + 2) = nvi_val;
504 nvi_val += nvi_val * pcts[3];
505 *out_ptr.add(i + 3) = nvi_val;
506 nvi_val += nvi_val * pcts[4];
507 *out_ptr.add(i + 4) = nvi_val;
508 nvi_val += nvi_val * pcts[5];
509 *out_ptr.add(i + 5) = nvi_val;
510 nvi_val += nvi_val * pcts[6];
511 *out_ptr.add(i + 6) = nvi_val;
512 nvi_val += nvi_val * pcts[7];
513 *out_ptr.add(i + 7) = nvi_val;
514
515 i += 8;
516 }
517
518 while i < len {
519 let c = *close_ptr.add(i);
520 let v = *vol_ptr.add(i);
521
522 if v < *vol_ptr.add(i - 1) {
523 let pct = (c - *close_ptr.add(i - 1)) / *close_ptr.add(i - 1);
524 nvi_val += nvi_val * pct;
525 }
526 *out_ptr.add(i) = nvi_val;
527 i += 1;
528 }
529}
530#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
531#[inline]
532pub unsafe fn nvi_avx512_short(close: &[f64], volume: &[f64], first: usize, out: &mut [f64]) {
533 nvi_avx512(close, volume, first, out)
534}
535#[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
536#[inline]
537pub unsafe fn nvi_avx512_long(close: &[f64], volume: &[f64], first: usize, out: &mut [f64]) {
538 nvi_avx512(close, volume, first, out)
539}
540
541#[inline(always)]
542pub fn nvi_batch_with_kernel(
543 close: &[f64],
544 volume: &[f64],
545 k: Kernel,
546) -> Result<NviBatchOutput, NviError> {
547 if close.is_empty() || volume.is_empty() {
548 return Err(NviError::EmptyInputData);
549 }
550 if close.len() != volume.len() {
551 return Err(NviError::MismatchedLength {
552 close_len: close.len(),
553 volume_len: volume.len(),
554 });
555 }
556
557 let cols = close.len();
558 let first = close
559 .iter()
560 .zip(volume)
561 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
562 .ok_or_else(|| {
563 if close.iter().all(|&c| c.is_nan()) {
564 NviError::AllCloseValuesNaN
565 } else {
566 NviError::AllVolumeValuesNaN
567 }
568 })?;
569 if cols - first < 2 {
570 return Err(NviError::NotEnoughValidData {
571 needed: 2,
572 valid: cols - first,
573 });
574 }
575
576 let mut buf_mu = make_uninit_matrix(1, cols);
577 init_matrix_prefixes(&mut buf_mu, cols, &[first]);
578
579 let mut guard = core::mem::ManuallyDrop::new(buf_mu);
580 let out: &mut [f64] =
581 unsafe { core::slice::from_raw_parts_mut(guard.as_mut_ptr() as *mut f64, guard.len()) };
582
583 let chosen = match k {
584 Kernel::Auto => detect_best_batch_kernel(),
585 other if other.is_batch() => other,
586 other => return Err(NviError::InvalidKernelForBatch(other)),
587 };
588 unsafe {
589 match chosen {
590 Kernel::Scalar | Kernel::ScalarBatch => nvi_row_scalar(close, volume, first, out),
591 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
592 Kernel::Avx2 | Kernel::Avx2Batch => nvi_row_scalar(close, volume, first, out),
593 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
594 Kernel::Avx512 | Kernel::Avx512Batch => nvi_row_scalar(close, volume, first, out),
595 _ => unreachable!(),
596 }
597 }
598
599 let values = unsafe {
600 Vec::from_raw_parts(
601 guard.as_mut_ptr() as *mut f64,
602 guard.len(),
603 guard.capacity(),
604 )
605 };
606 Ok(NviBatchOutput {
607 values,
608 rows: 1,
609 cols,
610 })
611}
612
613#[inline(always)]
614unsafe fn nvi_row_scalar(close: &[f64], volume: &[f64], first: usize, row_out_flat: &mut [f64]) {
615 let len = close.len();
616 let out = &mut row_out_flat[..len];
617 let mut nvi_val = 1000.0;
618 out[first] = nvi_val;
619
620 let mut prev_close = close[first];
621 let mut prev_volume = volume[first];
622
623 for i in (first + 1)..len {
624 if volume[i] < prev_volume {
625 let pct = (close[i] - prev_close) / prev_close;
626 nvi_val += nvi_val * pct;
627 }
628 out[i] = nvi_val;
629 prev_close = close[i];
630 prev_volume = volume[i];
631 }
632}
633
634#[cfg(test)]
635mod tests {
636 use super::*;
637 use crate::skip_if_unsupported;
638 use crate::utilities::data_loader::read_candles_from_csv;
639
640 fn check_nvi_partial_params(
641 test_name: &str,
642 kernel: Kernel,
643 ) -> Result<(), Box<dyn std::error::Error>> {
644 skip_if_unsupported!(kernel, test_name);
645 let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
646 let candles = read_candles_from_csv(file_path)?;
647 let input = NviInput::with_default_candles(&candles);
648 let output = nvi_with_kernel(&input, kernel)?;
649 assert_eq!(output.values.len(), candles.close.len());
650 Ok(())
651 }
652
653 fn check_nvi_accuracy(
654 test_name: &str,
655 kernel: Kernel,
656 ) -> Result<(), Box<dyn std::error::Error>> {
657 skip_if_unsupported!(kernel, test_name);
658 let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
659 let candles = read_candles_from_csv(file_path)?;
660 let input = NviInput::with_default_candles(&candles);
661 let result = nvi_with_kernel(&input, kernel)?;
662 let expected_last_five = [
663 154243.6925373456,
664 153973.11239019397,
665 153973.11239019397,
666 154275.63921207888,
667 154275.63921207888,
668 ];
669 let start = result.values.len().saturating_sub(5);
670 for (i, &val) in result.values[start..].iter().enumerate() {
671 let diff = (val - expected_last_five[i]).abs();
672 assert!(
673 diff < 1e-5,
674 "[{}] NVI {:?} mismatch at idx {}: got {}, expected {}",
675 test_name,
676 kernel,
677 i,
678 val,
679 expected_last_five[i]
680 );
681 }
682 Ok(())
683 }
684
685 fn check_nvi_empty_data(
686 test_name: &str,
687 kernel: Kernel,
688 ) -> Result<(), Box<dyn std::error::Error>> {
689 skip_if_unsupported!(kernel, test_name);
690 let close_data: [f64; 0] = [];
691 let volume_data: [f64; 0] = [];
692 let input = NviInput::from_slices(&close_data, &volume_data, NviParams);
693 let res = nvi_with_kernel(&input, kernel);
694 assert!(
695 res.is_err(),
696 "[{}] NVI should fail with empty data",
697 test_name
698 );
699 Ok(())
700 }
701
702 fn check_nvi_not_enough_valid_data(
703 test_name: &str,
704 kernel: Kernel,
705 ) -> Result<(), Box<dyn std::error::Error>> {
706 skip_if_unsupported!(kernel, test_name);
707 let close_data = [f64::NAN, 100.0];
708 let volume_data = [f64::NAN, 120.0];
709 let input = NviInput::from_slices(&close_data, &volume_data, NviParams);
710 let res = nvi_with_kernel(&input, kernel);
711 assert!(
712 res.is_err(),
713 "[{}] NVI should fail with not enough valid data",
714 test_name
715 );
716 Ok(())
717 }
718
719 fn check_nvi_streaming(
720 test_name: &str,
721 kernel: Kernel,
722 ) -> Result<(), Box<dyn std::error::Error>> {
723 skip_if_unsupported!(kernel, test_name);
724 let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
725 let candles = read_candles_from_csv(file_path)?;
726 let close = candles.select_candle_field("close")?;
727 let volume = candles.select_candle_field("volume")?;
728 let input = NviInput::from_slices(close, volume, NviParams);
729 let batch_output = nvi_with_kernel(&input, kernel)?.values;
730 let mut stream = NviStream::try_new()?;
731
732 let first_valid = close
733 .iter()
734 .zip(volume.iter())
735 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
736 .unwrap_or(0);
737
738 let mut stream_values = alloc_with_nan_prefix(close.len(), first_valid);
739
740 for (i, (&c, &v)) in close.iter().zip(volume.iter()).enumerate() {
741 if let Some(nvi_val) = stream.update(c, v) {
742 stream_values[i] = nvi_val;
743 }
744 }
745 assert_eq!(batch_output.len(), stream_values.len());
746 for (i, (&b, &s)) in batch_output.iter().zip(stream_values.iter()).enumerate() {
747 if b.is_nan() && s.is_nan() {
748 continue;
749 }
750 let diff = (b - s).abs();
751 assert!(
752 diff < 1e-9,
753 "[{}] NVI streaming mismatch at idx {}: batch={}, stream={}, diff={}",
754 test_name,
755 i,
756 b,
757 s,
758 diff
759 );
760 }
761 Ok(())
762 }
763
764 #[cfg(debug_assertions)]
765 fn check_nvi_no_poison(
766 test_name: &str,
767 kernel: Kernel,
768 ) -> Result<(), Box<dyn std::error::Error>> {
769 skip_if_unsupported!(kernel, test_name);
770
771 let file_path = "src/data/2018-09-01-2024-Bitfinex_Spot-4h.csv";
772 let candles = read_candles_from_csv(file_path)?;
773
774 let test_scenarios = vec![
775 ("default_candles", NviInput::with_default_candles(&candles)),
776 (
777 "close_source",
778 NviInput::from_candles(&candles, "close", NviParams),
779 ),
780 (
781 "high_source",
782 NviInput::from_candles(&candles, "high", NviParams),
783 ),
784 (
785 "low_source",
786 NviInput::from_candles(&candles, "low", NviParams),
787 ),
788 (
789 "open_source",
790 NviInput::from_candles(&candles, "open", NviParams),
791 ),
792 ];
793
794 for (scenario_idx, (scenario_name, input)) in test_scenarios.iter().enumerate() {
795 let output = nvi_with_kernel(input, kernel)?;
796
797 for (i, &val) in output.values.iter().enumerate() {
798 if val.is_nan() {
799 continue;
800 }
801
802 let bits = val.to_bits();
803
804 if bits == 0x11111111_11111111 {
805 panic!(
806 "[{}] Found alloc_with_nan_prefix poison value {} (0x{:016X}) at index {} \
807 with scenario: {} (scenario set {})",
808 test_name, val, bits, i, scenario_name, scenario_idx
809 );
810 }
811
812 if bits == 0x22222222_22222222 {
813 panic!(
814 "[{}] Found init_matrix_prefixes poison value {} (0x{:016X}) at index {} \
815 with scenario: {} (scenario set {})",
816 test_name, val, bits, i, scenario_name, scenario_idx
817 );
818 }
819
820 if bits == 0x33333333_33333333 {
821 panic!(
822 "[{}] Found make_uninit_matrix poison value {} (0x{:016X}) at index {} \
823 with scenario: {} (scenario set {})",
824 test_name, val, bits, i, scenario_name, scenario_idx
825 );
826 }
827 }
828 }
829
830 Ok(())
831 }
832
833 #[cfg(not(debug_assertions))]
834 fn check_nvi_no_poison(
835 _test_name: &str,
836 _kernel: Kernel,
837 ) -> Result<(), Box<dyn std::error::Error>> {
838 Ok(())
839 }
840
841 #[cfg(test)]
842 fn check_nvi_property(
843 test_name: &str,
844 kernel: Kernel,
845 ) -> Result<(), Box<dyn std::error::Error>> {
846 use proptest::prelude::*;
847 skip_if_unsupported!(kernel, test_name);
848
849 let strat = (50usize..=500)
850 .prop_flat_map(|len| {
851 (
852 prop::collection::vec(
853 prop::strategy::Union::new(vec![
854 (0.001f64..0.1f64).boxed(),
855 (10f64..10000f64).boxed(),
856 (1e6f64..1e8f64).boxed(),
857 ])
858 .prop_filter("finite", |x| x.is_finite()),
859 len,
860 ),
861 prop::collection::vec(
862 prop::strategy::Union::new(vec![
863 (100f64..1000f64).boxed(),
864 (1000f64..1e6f64).boxed(),
865 (1e6f64..1e9f64).boxed(),
866 ])
867 .prop_filter("finite", |x| x.is_finite()),
868 len,
869 ),
870 0usize..=7,
871 )
872 })
873 .prop_map(|(mut prices, mut volumes, scenario)| {
874 match scenario {
875 0 => {}
876 1 => {
877 let const_vol = volumes[0];
878 volumes.iter_mut().for_each(|v| *v = const_vol);
879 }
880 2 => {
881 volumes.sort_by(|a, b| b.partial_cmp(a).unwrap());
882 }
883 3 => {
884 volumes.sort_by(|a, b| a.partial_cmp(b).unwrap());
885 }
886 4 => {
887 for i in 0..volumes.len() {
888 volumes[i] = if i % 2 == 0 { 1000.0 } else { 500.0 };
889 }
890 }
891 5 => {
892 let const_price = prices[0];
893 prices.iter_mut().for_each(|p| *p = const_price);
894 }
895 6 => {
896 let start = prices[0];
897 let trend = 0.01f64;
898 for i in 0..prices.len() {
899 prices[i] = start * (1.0 + trend).powi(i as i32);
900 }
901 }
902 7 => {
903 let base = prices[0];
904 for i in 0..prices.len() {
905 prices[i] = base * (1.0 + 0.1 * ((i as f64 * 0.5).sin()));
906 }
907
908 for i in 0..volumes.len() {
909 volumes[i] *= (1.0 - (i as f64 / volumes.len() as f64) * 0.5);
910 }
911 }
912 _ => unreachable!(),
913 }
914 (prices, volumes, scenario)
915 });
916
917 proptest::test_runner::TestRunner::default()
918 .run(&strat, |(close_data, volume_data, scenario)| {
919 let input = NviInput::from_slices(&close_data, &volume_data, NviParams);
920
921 let NviOutput { values: out } = nvi_with_kernel(&input, kernel)?;
922
923 let NviOutput { values: ref_out } = nvi_with_kernel(&input, Kernel::Scalar)?;
924
925 let first_valid = close_data
926 .iter()
927 .zip(volume_data.iter())
928 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
929 .unwrap_or(close_data.len());
930
931 if first_valid >= close_data.len() {
932 return Ok(());
933 }
934
935 prop_assert!(
936 (out[first_valid] - 1000.0).abs() < 1e-9,
937 "NVI should start at 1000.0, got {} at index {} (scenario {})",
938 out[first_valid],
939 first_valid,
940 scenario
941 );
942
943 let mut prev_nvi = 1000.0;
944 let mut prev_close = close_data[first_valid];
945 let mut prev_volume = volume_data[first_valid];
946
947 for i in (first_valid + 1)..close_data.len() {
948 let curr_close = close_data[i];
949 let curr_volume = volume_data[i];
950 let curr_nvi = out[i];
951
952 if curr_volume < prev_volume {
953 let expected_pct = (curr_close - prev_close) / prev_close;
954 let expected_nvi = prev_nvi + prev_nvi * expected_pct;
955
956 prop_assert!(
957 (curr_nvi - expected_nvi).abs() < 1e-9 ||
958 (curr_nvi - expected_nvi).abs() / expected_nvi.abs() < 1e-9,
959 "NVI calculation error at index {} (scenario {}): expected {}, got {}, \
960 prev_nvi={}, pct_change={}, volume {} -> {}",
961 i, scenario, expected_nvi, curr_nvi, prev_nvi, expected_pct,
962 prev_volume, curr_volume
963 );
964 } else {
965 prop_assert!(
966 (curr_nvi - prev_nvi).abs() < 1e-9,
967 "NVI should not change when volume doesn't decrease at index {} (scenario {}): \
968 prev_nvi={}, curr_nvi={}, volume {} -> {}",
969 i, scenario, prev_nvi, curr_nvi, prev_volume, curr_volume
970 );
971 }
972
973 prev_nvi = curr_nvi;
974 prev_close = curr_close;
975 prev_volume = curr_volume;
976 }
977
978 for i in first_valid..close_data.len() {
979 let y = out[i];
980 let r = ref_out[i];
981
982 if !y.is_finite() || !r.is_finite() {
983 prop_assert!(
984 y.to_bits() == r.to_bits(),
985 "Kernel finite/NaN mismatch at index {} (scenario {}): {} vs {}",
986 i,
987 scenario,
988 y,
989 r
990 );
991 } else {
992 let ulp_diff = y.to_bits().abs_diff(r.to_bits());
993 prop_assert!(
994 (y - r).abs() <= 1e-9 || ulp_diff <= 4,
995 "Kernel mismatch at index {} (scenario {}): {} vs {} (ULP={})",
996 i,
997 scenario,
998 y,
999 r,
1000 ulp_diff
1001 );
1002 }
1003 }
1004
1005 match scenario {
1006 1 => {
1007 for i in (first_valid + 1)..out.len() {
1008 prop_assert!(
1009 (out[i] - 1000.0).abs() < 1e-9,
1010 "NVI should stay at 1000.0 with constant volume, got {} at index {}",
1011 out[i], i
1012 );
1013 }
1014 }
1015 3 => {
1016 for i in (first_valid + 1)..out.len() {
1017 prop_assert!(
1018 (out[i] - 1000.0).abs() < 1e-9,
1019 "NVI should stay at 1000.0 with always increasing volume, got {} at index {}",
1020 out[i], i
1021 );
1022 }
1023 }
1024 5 => {
1025 if first_valid + 1 < out.len() {
1026 let mut expected_nvi = out[first_valid];
1027 for i in (first_valid + 1)..out.len() {
1028 prop_assert!(
1029 (out[i] - expected_nvi).abs() < 1e-9,
1030 "NVI should stay constant at {} with constant prices, got {} at index {}",
1031 expected_nvi, out[i], i
1032 );
1033 }
1034 }
1035 }
1036 _ => {}
1037 }
1038
1039 let mut stream = NviStream::try_new()?;
1040 for i in 0..close_data.len() {
1041 if let Some(stream_val) = stream.update(close_data[i], volume_data[i]) {
1042 let batch_val = out[i];
1043 if !batch_val.is_nan() {
1044 prop_assert!(
1045 (stream_val - batch_val).abs() < 1e-9,
1046 "Streaming mismatch at index {} (scenario {}): stream={}, batch={}",
1047 i,
1048 scenario,
1049 stream_val,
1050 batch_val
1051 );
1052 }
1053 }
1054 }
1055
1056 Ok(())
1057 })
1058 .unwrap();
1059
1060 Ok(())
1061 }
1062
1063 macro_rules! generate_all_nvi_tests {
1064 ($($test_fn:ident),*) => {
1065 paste::paste! {
1066 $( #[test] fn [<$test_fn _scalar_f64>]() { let _ = $test_fn(stringify!([<$test_fn _scalar_f64>]), Kernel::Scalar); } )*
1067 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
1068 $( #[test] fn [<$test_fn _avx2_f64>]() { let _ = $test_fn(stringify!([<$test_fn _avx2_f64>]), Kernel::Avx2); } )*
1069 #[cfg(all(feature = "nightly-avx", target_arch = "x86_64"))]
1070 $( #[test] fn [<$test_fn _avx512_f64>]() { let _ = $test_fn(stringify!([<$test_fn _avx512_f64>]), Kernel::Avx512); } )*
1071 }
1072 }
1073 }
1074
1075 generate_all_nvi_tests!(
1076 check_nvi_partial_params,
1077 check_nvi_accuracy,
1078 check_nvi_empty_data,
1079 check_nvi_not_enough_valid_data,
1080 check_nvi_streaming,
1081 check_nvi_no_poison
1082 );
1083
1084 #[cfg(test)]
1085 generate_all_nvi_tests!(check_nvi_property);
1086
1087 #[test]
1088 fn test_nvi_into_matches_api() -> Result<(), Box<dyn std::error::Error>> {
1089 let len = 256usize;
1090 let mut close = vec![f64::NAN; len];
1091 let mut volume = vec![f64::NAN; len];
1092
1093 for i in 5..len {
1094 let t = (i - 5) as f64;
1095
1096 close[i] = 100.0 + 0.05 * t + (0.01 * t).sin();
1097
1098 volume[i] = 2000.0 + ((i as i64 % 7) as f64 - 3.0) * 40.0;
1099 }
1100
1101 let input = NviInput::from_slices(&close, &volume, NviParams);
1102
1103 let baseline = nvi(&input)?.values;
1104
1105 let mut out = vec![0.0; len];
1106 #[cfg(not(all(target_arch = "wasm32", feature = "wasm")))]
1107 {
1108 nvi_into(&input, &mut out)?;
1109 }
1110 #[cfg(all(target_arch = "wasm32", feature = "wasm"))]
1111 {
1112 nvi_into_slice(&mut out, &close, &volume, Kernel::Auto)?;
1113 }
1114
1115 assert_eq!(baseline.len(), out.len());
1116 for (i, (&a, &b)) in baseline.iter().zip(out.iter()).enumerate() {
1117 let equal = (a.is_nan() && b.is_nan()) || (a - b).abs() <= 1e-12;
1118 assert!(
1119 equal,
1120 "nvi_into parity mismatch at index {}: {} vs {}",
1121 i, a, b
1122 );
1123 }
1124 Ok(())
1125 }
1126}
1127
1128#[cfg(feature = "python")]
1129#[pyclass(name = "NviStream")]
1130pub struct NviStreamPy {
1131 stream: NviStream,
1132}
1133
1134#[cfg(feature = "python")]
1135#[pymethods]
1136impl NviStreamPy {
1137 #[new]
1138 fn new() -> PyResult<Self> {
1139 let stream = NviStream::try_new().map_err(|e| PyValueError::new_err(e.to_string()))?;
1140 Ok(NviStreamPy { stream })
1141 }
1142
1143 fn update(&mut self, close: f64, volume: f64) -> Option<f64> {
1144 self.stream.update(close, volume)
1145 }
1146}
1147
1148#[cfg(feature = "python")]
1149#[pyfunction(name = "nvi")]
1150#[pyo3(signature = (close, volume, kernel=None))]
1151pub fn nvi_py<'py>(
1152 py: Python<'py>,
1153 close: PyReadonlyArray1<'py, f64>,
1154 volume: PyReadonlyArray1<'py, f64>,
1155 kernel: Option<&str>,
1156) -> PyResult<Bound<'py, PyArray1<f64>>> {
1157 let close_slice = close.as_slice()?;
1158 let volume_slice = volume.as_slice()?;
1159 let kern = validate_kernel(kernel, false)?;
1160
1161 let input = NviInput::from_slices(close_slice, volume_slice, NviParams);
1162
1163 let result_vec: Vec<f64> = py
1164 .allow_threads(|| nvi_with_kernel(&input, kern).map(|o| o.values))
1165 .map_err(|e| PyValueError::new_err(e.to_string()))?;
1166
1167 Ok(result_vec.into_pyarray(py))
1168}
1169
1170#[cfg(feature = "python")]
1171#[pyfunction(name = "nvi_batch")]
1172#[pyo3(signature = (close, volume, kernel=None))]
1173pub fn nvi_batch_py<'py>(
1174 py: Python<'py>,
1175 close: PyReadonlyArray1<'py, f64>,
1176 volume: PyReadonlyArray1<'py, f64>,
1177 kernel: Option<&str>,
1178) -> PyResult<Bound<'py, pyo3::types::PyDict>> {
1179 use numpy::{IntoPyArray, PyArray1, PyArrayMethods};
1180 use pyo3::types::PyDict;
1181
1182 let close_slice = close.as_slice()?;
1183 let volume_slice = volume.as_slice()?;
1184 let kern = validate_kernel(kernel, true)?;
1185
1186 let rows = 1usize;
1187 let cols = close_slice.len();
1188 let out_arr = unsafe { PyArray1::<f64>::new(py, [rows * cols], false) };
1189 let out_slice = unsafe { out_arr.as_slice_mut()? };
1190
1191 py.allow_threads(|| -> Result<(), NviError> {
1192 if close_slice.len() != volume_slice.len() {
1193 return Err(NviError::MismatchedLength {
1194 close_len: close_slice.len(),
1195 volume_len: volume_slice.len(),
1196 });
1197 }
1198 let first = close_slice
1199 .iter()
1200 .zip(volume_slice)
1201 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
1202 .ok_or_else(|| {
1203 if close_slice.iter().all(|&c| c.is_nan()) {
1204 NviError::AllCloseValuesNaN
1205 } else {
1206 NviError::AllVolumeValuesNaN
1207 }
1208 })?;
1209 if cols - first < 2 {
1210 return Err(NviError::NotEnoughValidData {
1211 needed: 2,
1212 valid: cols - first,
1213 });
1214 }
1215
1216 for v in &mut out_slice[..first] {
1217 *v = f64::NAN;
1218 }
1219
1220 unsafe { nvi_row_scalar(close_slice, volume_slice, first, out_slice) };
1221 Ok(())
1222 })
1223 .map_err(|e| PyValueError::new_err(e.to_string()))?;
1224
1225 let d = PyDict::new(py);
1226 d.set_item("values", out_arr.reshape((rows, cols))?)?;
1227 d.set_item("rows", rows)?;
1228 d.set_item("cols", cols)?;
1229 Ok(d)
1230}
1231
1232#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
1233#[wasm_bindgen]
1234pub fn nvi_js(close: &[f64], volume: &[f64]) -> Result<Vec<f64>, JsValue> {
1235 let mut output = vec![0.0; close.len()];
1236
1237 nvi_into_slice(&mut output, close, volume, Kernel::Auto)
1238 .map_err(|e| JsValue::from_str(&e.to_string()))?;
1239
1240 Ok(output)
1241}
1242
1243#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
1244#[wasm_bindgen]
1245pub fn nvi_into(
1246 close_ptr: *const f64,
1247 volume_ptr: *const f64,
1248 out_ptr: *mut f64,
1249 len: usize,
1250) -> Result<(), JsValue> {
1251 if close_ptr.is_null() || volume_ptr.is_null() || out_ptr.is_null() {
1252 return Err(JsValue::from_str("Null pointer provided"));
1253 }
1254
1255 unsafe {
1256 let close = std::slice::from_raw_parts(close_ptr, len);
1257 let volume = std::slice::from_raw_parts(volume_ptr, len);
1258
1259 if close_ptr == out_ptr as *const f64 || volume_ptr == out_ptr as *const f64 {
1260 let mut temp = vec![0.0; len];
1261 nvi_into_slice(&mut temp, close, volume, Kernel::Auto)
1262 .map_err(|e| JsValue::from_str(&e.to_string()))?;
1263 let out = std::slice::from_raw_parts_mut(out_ptr, len);
1264 out.copy_from_slice(&temp);
1265 } else {
1266 let out = std::slice::from_raw_parts_mut(out_ptr, len);
1267 nvi_into_slice(out, close, volume, Kernel::Auto)
1268 .map_err(|e| JsValue::from_str(&e.to_string()))?;
1269 }
1270 Ok(())
1271 }
1272}
1273
1274#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
1275#[wasm_bindgen]
1276pub fn nvi_alloc(len: usize) -> *mut f64 {
1277 let mut vec = Vec::<f64>::with_capacity(len);
1278 let ptr = vec.as_mut_ptr();
1279 std::mem::forget(vec);
1280 ptr
1281}
1282
1283#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
1284#[wasm_bindgen]
1285pub fn nvi_free(ptr: *mut f64, len: usize) {
1286 if !ptr.is_null() {
1287 unsafe {
1288 let _ = Vec::from_raw_parts(ptr, len, len);
1289 }
1290 }
1291}
1292
1293#[cfg(all(target_arch = "wasm32", feature = "wasm"))]
1294#[wasm_bindgen]
1295pub fn nvi_batch_into(
1296 close_ptr: *const f64,
1297 volume_ptr: *const f64,
1298 out_ptr: *mut f64,
1299 len: usize,
1300) -> Result<usize, JsValue> {
1301 if close_ptr.is_null() || volume_ptr.is_null() || out_ptr.is_null() {
1302 return Err(JsValue::from_str("Null pointer provided"));
1303 }
1304 unsafe {
1305 let close = std::slice::from_raw_parts(close_ptr, len);
1306 let volume = std::slice::from_raw_parts(volume_ptr, len);
1307 let out = std::slice::from_raw_parts_mut(out_ptr, len);
1308
1309 if close.len() != volume.len() {
1310 return Err(JsValue::from_str("Length mismatch"));
1311 }
1312 let first = close
1313 .iter()
1314 .zip(volume)
1315 .position(|(&c, &v)| !c.is_nan() && !v.is_nan())
1316 .ok_or_else(|| JsValue::from_str("All values NaN in one or both inputs"))?;
1317 if len - first < 2 {
1318 return Err(JsValue::from_str("Not enough valid data"));
1319 }
1320
1321 for v in &mut out[..first] {
1322 *v = f64::NAN;
1323 }
1324 nvi_row_scalar(close, volume, first, out);
1325 Ok(1)
1326 }
1327}
1328
1329#[cfg(all(feature = "python", feature = "cuda"))]
1330use crate::cuda::cuda_available;
1331#[cfg(all(feature = "python", feature = "cuda"))]
1332use crate::cuda::CudaNvi;
1333#[cfg(all(feature = "python", feature = "cuda"))]
1334use crate::indicators::moving_averages::alma::DeviceArrayF32Py;
1335
1336#[cfg(all(feature = "python", feature = "cuda"))]
1337#[pyfunction(name = "nvi_cuda_batch_dev")]
1338#[pyo3(signature = (close, volume, device_id=0))]
1339pub fn nvi_cuda_batch_dev_py(
1340 py: Python<'_>,
1341 close: PyReadonlyArray1<'_, f32>,
1342 volume: PyReadonlyArray1<'_, f32>,
1343 device_id: usize,
1344) -> PyResult<DeviceArrayF32Py> {
1345 if !cuda_available() {
1346 return Err(PyValueError::new_err("CUDA not available"));
1347 }
1348 let close_slice = close.as_slice()?;
1349 let volume_slice = volume.as_slice()?;
1350 if close_slice.len() != volume_slice.len() {
1351 return Err(PyValueError::new_err("mismatched input lengths"));
1352 }
1353 let (inner, ctx, dev_id) = py.allow_threads(|| {
1354 let cuda = CudaNvi::new(device_id).map_err(|e| PyValueError::new_err(e.to_string()))?;
1355 let ctx = cuda.context_arc();
1356 let dev_id = cuda.device_id();
1357 let arr = cuda
1358 .nvi_batch_dev(close_slice, volume_slice)
1359 .map_err(|e| PyValueError::new_err(e.to_string()))?;
1360 Ok::<_, PyErr>((arr, ctx, dev_id))
1361 })?;
1362 Ok(DeviceArrayF32Py {
1363 inner,
1364 _ctx: Some(ctx),
1365 device_id: Some(dev_id),
1366 })
1367}
1368
1369#[cfg(all(feature = "python", feature = "cuda"))]
1370#[pyfunction(name = "nvi_cuda_many_series_one_param_dev")]
1371#[pyo3(signature = (close_tm, volume_tm, cols, rows, device_id=0))]
1372pub fn nvi_cuda_many_series_one_param_dev_py(
1373 py: Python<'_>,
1374 close_tm: PyReadonlyArray1<'_, f32>,
1375 volume_tm: PyReadonlyArray1<'_, f32>,
1376 cols: usize,
1377 rows: usize,
1378 device_id: usize,
1379) -> PyResult<DeviceArrayF32Py> {
1380 if !cuda_available() {
1381 return Err(PyValueError::new_err("CUDA not available"));
1382 }
1383 let close_slice = close_tm.as_slice()?;
1384 let volume_slice = volume_tm.as_slice()?;
1385 let (inner, ctx, dev_id) = py.allow_threads(|| {
1386 let cuda = CudaNvi::new(device_id).map_err(|e| PyValueError::new_err(e.to_string()))?;
1387 let ctx = cuda.context_arc();
1388 let dev_id = cuda.device_id();
1389 let arr = cuda
1390 .nvi_many_series_one_param_time_major_dev(close_slice, volume_slice, cols, rows)
1391 .map_err(|e| PyValueError::new_err(e.to_string()))?;
1392 Ok::<_, PyErr>((arr, ctx, dev_id))
1393 })?;
1394 Ok(DeviceArrayF32Py {
1395 inner,
1396 _ctx: Some(ctx),
1397 device_id: Some(dev_id),
1398 })
1399}