refactor(cuda): replace Candle Tensor with cudarc CudaSlice in signal_adapter
Eliminate all 26 Candle Tensor references from signal_adapter.rs. Three functions (ppo_to_exposure_scores, signal_to_action_scores, tft_quantile_to_signal) now take CudaSlice<f32> + CudaStream and return CudaSlice<f32>, backed by three fused CUDA kernels in signal_adapter_kernel.cu. Dead code evaluate_supervised_gpu_backtest removed (zero callers). Added cuda_f32_to_tensor helper to gpu_action_selector for DtoD copy back to Candle Tensor at API boundaries (PPO hyperopt adapter, evaluate_baseline example). Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
This commit is contained in:
jgrusewski
@@ -1506,6 +1506,8 @@ fn evaluate_supervised_fold_gpu(
|
||||
) -> Result<Vec<ml::cuda_pipeline::gpu_backtest_evaluator::WindowMetrics>> {
|
||||
use ml::cuda_pipeline::gpu_backtest_evaluator::{GpuBacktestConfig, GpuBacktestEvaluator};
|
||||
use ml::cuda_pipeline::signal_adapter::{signal_to_action_scores, tft_quantile_to_signal};
|
||||
use ml::cuda_pipeline::gpu_action_selector::cuda_f32_to_tensor;
|
||||
use candle_core::cuda_backend::cudarc::driver::CudaSlice;
|
||||
use std::cell::RefCell;
|
||||
|
||||
let device = candle_core::Device::cuda_if_available(0)
|
||||
@@ -1604,32 +1606,62 @@ fn evaluate_supervised_fold_gpu(
|
||||
.map_err(|e| ml::MLError::ModelError(format!("borrow_mut: {e}")))?;
|
||||
|
||||
let raw_output = model_ref.forward(&market_input)?;
|
||||
let batch = raw_output.dims().first().copied().unwrap_or(1);
|
||||
|
||||
// Convert raw output → scalar signal → [batch, 5] action scores
|
||||
let signal = if is_tft {
|
||||
// TFT outputs [batch, horizon, num_quantiles] — extract median
|
||||
tft_quantile_to_signal(&raw_output)?
|
||||
let cuda_dev = match &device {
|
||||
candle_core::Device::Cuda(d) => d,
|
||||
_ => return Err(ml::MLError::ModelError("device not CUDA".into())),
|
||||
};
|
||||
let stream = cuda_dev.cuda_stream();
|
||||
|
||||
// Convert raw output -> scalar signal CudaSlice -> [batch, 5] action scores
|
||||
if is_tft {
|
||||
// TFT outputs [batch, horizon, num_quantiles] -- extract median via CUDA kernel
|
||||
let dims = raw_output.dims();
|
||||
let horizon = dims.get(1).copied().unwrap_or(1);
|
||||
let num_q = dims.get(2).copied().unwrap_or(3);
|
||||
let (guard, _layout) = raw_output.storage_and_layout();
|
||||
let q_slice: &CudaSlice<f32> = match &*guard {
|
||||
candle_core::Storage::Cuda(cs) => cs.as_cuda_slice()
|
||||
.map_err(|e| ml::MLError::ModelError(format!("tft as_cuda_slice: {e}")))?,
|
||||
_ => return Err(ml::MLError::ModelError("TFT output not on CUDA".into())),
|
||||
};
|
||||
let signal_slice = tft_quantile_to_signal(q_slice, batch, horizon, num_q, &stream)?;
|
||||
let scores_slice = signal_to_action_scores(
|
||||
&signal_slice, batch, signal_high, signal_low, &stream,
|
||||
)?;
|
||||
drop(guard);
|
||||
cuda_f32_to_tensor(&scores_slice, &[batch, 5], &device)
|
||||
} else {
|
||||
// Other supervised models output [batch, 1] or [batch] scalar
|
||||
if raw_output.dims().len() == 2 {
|
||||
let squeezed = if raw_output.dims().len() == 2 {
|
||||
let last_dim = raw_output.dims().get(1).copied().unwrap_or(0);
|
||||
if last_dim == 1 {
|
||||
raw_output
|
||||
.squeeze(1)
|
||||
.map_err(|e| ml::MLError::ModelError(format!("squeeze: {e}")))?
|
||||
} else {
|
||||
// For sequence models like Mamba2 that output [batch, seq, 1],
|
||||
// take the last time step
|
||||
raw_output
|
||||
.flatten_all()
|
||||
.map_err(|e| ml::MLError::ModelError(format!("flatten: {e}")))?
|
||||
}
|
||||
} else {
|
||||
raw_output
|
||||
}
|
||||
};
|
||||
};
|
||||
|
||||
signal_to_action_scores(&signal, signal_high, signal_low)
|
||||
// Extract CudaSlice, run kernel while guard is alive, wrap result
|
||||
let (guard, _layout) = squeezed.storage_and_layout();
|
||||
let s_slice: &CudaSlice<f32> = match &*guard {
|
||||
candle_core::Storage::Cuda(cs) => cs.as_cuda_slice()
|
||||
.map_err(|e| ml::MLError::ModelError(format!("signal as_cuda_slice: {e}")))?,
|
||||
_ => return Err(ml::MLError::ModelError("signal not on CUDA".into())),
|
||||
};
|
||||
let scores_slice = signal_to_action_scores(
|
||||
s_slice, batch, signal_high, signal_low, &stream,
|
||||
)?;
|
||||
drop(guard);
|
||||
cuda_f32_to_tensor(&scores_slice, &[batch, 5], &device)
|
||||
}
|
||||
},
|
||||
3, // portfolio_dim
|
||||
&device,
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,42 +1,106 @@
|
||||
#![allow(unsafe_code)] // Required for CUDA kernel launch
|
||||
|
||||
//! Signal adapter utilities for converting between model outputs and action scores.
|
||||
//!
|
||||
//! Pure Candle tensor ops — works on any device (CPU, CUDA, Metal).
|
||||
//! Four functions bridge the gap between model predictions and the 5-action
|
||||
//! Pure cudarc `CudaSlice<f32>` operations -- zero Candle Tensor usage.
|
||||
//! Three CUDA kernels bridge the gap between model predictions and the 5-action
|
||||
//! DQN exposure space (Short100, Short50, Flat, Long50, Long100).
|
||||
|
||||
use candle_core::{DType, Tensor};
|
||||
use candle_core::cuda_backend::cudarc;
|
||||
use cudarc::driver::{CudaContext, CudaFunction, CudaSlice, CudaStream, LaunchConfig, PushKernelArg};
|
||||
use cudarc::nvrtc::Ptx;
|
||||
use std::sync::OnceLock;
|
||||
|
||||
use crate::MLError;
|
||||
|
||||
// ── Kernel compilation ──────────────────────────────────────────────────
|
||||
|
||||
static SIGNAL_ADAPTER_PTX: OnceLock<Result<Ptx, String>> = OnceLock::new();
|
||||
|
||||
fn compile_signal_adapter_ptx(context: &CudaContext) -> Result<Ptx, String> {
|
||||
let kernel_src = include_str!("signal_adapter_kernel.cu");
|
||||
crate::cuda_pipeline::compile_ptx_for_device(kernel_src, context)
|
||||
}
|
||||
|
||||
/// Lazily compiled kernel handle set for signal adapter operations.
|
||||
struct KernelSet {
|
||||
ppo_exposure: CudaFunction,
|
||||
signal_to_action: CudaFunction,
|
||||
tft_extract: CudaFunction,
|
||||
}
|
||||
|
||||
fn load_kernels(context: &CudaContext) -> Result<KernelSet, MLError> {
|
||||
let ptx = SIGNAL_ADAPTER_PTX
|
||||
.get_or_init(|| compile_signal_adapter_ptx(context))
|
||||
.as_ref()
|
||||
.map_err(|e| MLError::ModelError(format!("signal_adapter PTX: {e}")))?;
|
||||
let module = context
|
||||
.load_module(ptx.clone())
|
||||
.map_err(|e| MLError::ModelError(format!("signal_adapter module load: {e}")))?;
|
||||
|
||||
let ppo_exposure = module
|
||||
.load_function("ppo_to_exposure_scores_kernel")
|
||||
.map_err(|e| MLError::ModelError(format!("ppo_to_exposure_scores_kernel load: {e}")))?;
|
||||
let signal_to_action = module
|
||||
.load_function("signal_to_action_scores_kernel")
|
||||
.map_err(|e| MLError::ModelError(format!("signal_to_action_scores_kernel load: {e}")))?;
|
||||
let tft_extract = module
|
||||
.load_function("tft_quantile_extract_kernel")
|
||||
.map_err(|e| MLError::ModelError(format!("tft_quantile_extract_kernel load: {e}")))?;
|
||||
|
||||
Ok(KernelSet { ppo_exposure, signal_to_action, tft_extract })
|
||||
}
|
||||
|
||||
// ── Public API ──────────────────────────────────────────────────────────
|
||||
|
||||
/// Aggregate PPO 45-action softmax probabilities into 5 exposure scores.
|
||||
///
|
||||
/// PPO uses a 45-action factored space (5 exposure x 3 order x 3 urgency).
|
||||
/// This collapses the order/urgency dimensions to produce a `[batch, 5]`
|
||||
/// tensor whose argmax gives the dominant exposure bucket.
|
||||
/// This collapses the order/urgency dimensions to produce a `[batch * 5]`
|
||||
/// `CudaSlice` whose per-row argmax gives the dominant exposure bucket.
|
||||
///
|
||||
/// # Arguments
|
||||
/// * `probs` - GPU-resident `[batch * 45]` softmax probabilities
|
||||
/// * `batch` - number of rows
|
||||
/// * `stream` - CUDA stream for kernel launch and allocation
|
||||
///
|
||||
/// # Returns
|
||||
/// GPU-resident `CudaSlice<f32>` of shape `[batch * 5]`.
|
||||
///
|
||||
/// # Errors
|
||||
/// Returns `MLError::ModelError` if the last dimension is not 45.
|
||||
pub fn ppo_to_exposure_scores(probs: &Tensor) -> Result<Tensor, MLError> {
|
||||
let shape = probs.dims();
|
||||
let ndim = shape.len();
|
||||
if ndim < 2 {
|
||||
return Err(MLError::ModelError(format!(
|
||||
"ppo_to_exposure_scores: expected >= 2 dims, got {ndim}"
|
||||
)));
|
||||
/// Returns `MLError::ModelError` on kernel compilation or launch failure.
|
||||
pub fn ppo_to_exposure_scores(
|
||||
probs: &CudaSlice<f32>,
|
||||
batch: usize,
|
||||
stream: &CudaStream,
|
||||
) -> Result<CudaSlice<f32>, MLError> {
|
||||
let context = stream.context();
|
||||
let kernels = load_kernels(&context)?;
|
||||
|
||||
let mut out = stream
|
||||
.alloc_zeros::<f32>(batch * 5)
|
||||
.map_err(|e| MLError::ModelError(format!("alloc exposure scores: {e}")))?;
|
||||
|
||||
let batch_i32 = batch as i32;
|
||||
let grid = ((batch as u32 + 255) / 256).max(1);
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (grid, 1, 1),
|
||||
block_dim: (256, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
|
||||
// Safety: argument order matches ppo_to_exposure_scores_kernel(probs, out, batch)
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernels.ppo_exposure)
|
||||
.arg(probs)
|
||||
.arg(&mut out)
|
||||
.arg(&batch_i32)
|
||||
.launch(cfg)
|
||||
.map_err(|e| MLError::ModelError(format!("ppo_to_exposure_scores launch: {e}")))?;
|
||||
}
|
||||
let last = shape.get(ndim - 1).copied().unwrap_or(0);
|
||||
if last != 45 {
|
||||
return Err(MLError::ModelError(format!(
|
||||
"ppo_to_exposure_scores: last dim must be 45, got {last}"
|
||||
)));
|
||||
}
|
||||
let batch = shape.get(0).copied().unwrap_or(1);
|
||||
let reshaped = probs
|
||||
.reshape(&[batch, 5, 9])
|
||||
.map_err(|e| MLError::ModelError(format!("reshape to [B,5,9]: {e}")))?;
|
||||
let summed = reshaped
|
||||
.sum(2)
|
||||
.map_err(|e| MLError::ModelError(format!("sum over urgency*order dim: {e}")))?;
|
||||
Ok(summed)
|
||||
|
||||
Ok(out)
|
||||
}
|
||||
|
||||
/// Map scalar return predictions (bps) to 5-action one-hot-ish scores.
|
||||
@@ -49,228 +113,124 @@ pub fn ppo_to_exposure_scores(probs: &Tensor) -> Result<Tensor, MLError> {
|
||||
/// - `+low < pred <= +high` => Long50 (action 3)
|
||||
/// - `pred > +high` => Long100 (action 4)
|
||||
///
|
||||
/// # Arguments
|
||||
/// * `predictions` - GPU-resident `[batch]` scalar signals
|
||||
/// * `batch` - number of elements
|
||||
/// * `high_threshold_bps` - strong signal threshold
|
||||
/// * `low_threshold_bps` - mild signal threshold
|
||||
/// * `stream` - CUDA stream for kernel launch and allocation
|
||||
///
|
||||
/// # Returns
|
||||
/// GPU-resident `CudaSlice<f32>` of shape `[batch * 5]` (one-hot rows).
|
||||
///
|
||||
/// # Errors
|
||||
/// Returns `MLError::ModelError` on shape or tensor-op failures.
|
||||
/// Returns `MLError::ModelError` on kernel compilation or launch failure.
|
||||
pub fn signal_to_action_scores(
|
||||
predictions: &Tensor,
|
||||
predictions: &CudaSlice<f32>,
|
||||
batch: usize,
|
||||
high_threshold_bps: f32,
|
||||
low_threshold_bps: f32,
|
||||
) -> Result<Tensor, MLError> {
|
||||
// Flatten to [batch]
|
||||
let pred = if predictions.dims().len() == 2 {
|
||||
predictions
|
||||
.squeeze(1)
|
||||
.map_err(|e| MLError::ModelError(format!("squeeze predictions: {e}")))?
|
||||
} else {
|
||||
predictions.clone()
|
||||
stream: &CudaStream,
|
||||
) -> Result<CudaSlice<f32>, MLError> {
|
||||
let context = stream.context();
|
||||
let kernels = load_kernels(&context)?;
|
||||
|
||||
let mut out = stream
|
||||
.alloc_zeros::<f32>(batch * 5)
|
||||
.map_err(|e| MLError::ModelError(format!("alloc action scores: {e}")))?;
|
||||
|
||||
let batch_i32 = batch as i32;
|
||||
let grid = ((batch as u32 + 255) / 256).max(1);
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (grid, 1, 1),
|
||||
block_dim: (256, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
|
||||
let device = pred.device();
|
||||
let batch = pred.dims().first().copied().unwrap_or(1);
|
||||
// Safety: argument order matches signal_to_action_scores_kernel(
|
||||
// predictions, out_scores, high_threshold_bps, low_threshold_bps, batch)
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernels.signal_to_action)
|
||||
.arg(predictions)
|
||||
.arg(&mut out)
|
||||
.arg(&high_threshold_bps)
|
||||
.arg(&low_threshold_bps)
|
||||
.arg(&batch_i32)
|
||||
.launch(cfg)
|
||||
.map_err(|e| MLError::ModelError(format!("signal_to_action_scores launch: {e}")))?;
|
||||
}
|
||||
|
||||
let high_pos = Tensor::new(&[high_threshold_bps], device)
|
||||
.map_err(|e| MLError::ModelError(format!("scalar high_pos: {e}")))?
|
||||
.broadcast_as(&[batch])
|
||||
.map_err(|e| MLError::ModelError(format!("broadcast high_pos: {e}")))?;
|
||||
let high_neg = Tensor::new(&[-high_threshold_bps], device)
|
||||
.map_err(|e| MLError::ModelError(format!("scalar high_neg: {e}")))?
|
||||
.broadcast_as(&[batch])
|
||||
.map_err(|e| MLError::ModelError(format!("broadcast high_neg: {e}")))?;
|
||||
let low_pos = Tensor::new(&[low_threshold_bps], device)
|
||||
.map_err(|e| MLError::ModelError(format!("scalar low_pos: {e}")))?
|
||||
.broadcast_as(&[batch])
|
||||
.map_err(|e| MLError::ModelError(format!("broadcast low_pos: {e}")))?;
|
||||
let low_neg = Tensor::new(&[-low_threshold_bps], device)
|
||||
.map_err(|e| MLError::ModelError(format!("scalar low_neg: {e}")))?
|
||||
.broadcast_as(&[batch])
|
||||
.map_err(|e| MLError::ModelError(format!("broadcast low_neg: {e}")))?;
|
||||
|
||||
// Short100: pred < -high
|
||||
let short100 = pred
|
||||
.lt(&high_neg)
|
||||
.map_err(|e| MLError::ModelError(format!("lt high_neg: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast short100: {e}")))?;
|
||||
|
||||
// Short50: pred >= -high AND pred < -low
|
||||
let ge_neg_high = pred
|
||||
.ge(&high_neg)
|
||||
.map_err(|e| MLError::ModelError(format!("ge high_neg: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast ge_neg_high: {e}")))?;
|
||||
let lt_neg_low = pred
|
||||
.lt(&low_neg)
|
||||
.map_err(|e| MLError::ModelError(format!("lt low_neg: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast lt_neg_low: {e}")))?;
|
||||
let short50 = ge_neg_high
|
||||
.mul(<_neg_low)
|
||||
.map_err(|e| MLError::ModelError(format!("mul short50: {e}")))?;
|
||||
|
||||
// Flat: pred >= -low AND pred <= +low
|
||||
let ge_neg_low = pred
|
||||
.ge(&low_neg)
|
||||
.map_err(|e| MLError::ModelError(format!("ge low_neg: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast ge_neg_low: {e}")))?;
|
||||
let le_pos_low = pred
|
||||
.le(&low_pos)
|
||||
.map_err(|e| MLError::ModelError(format!("le low_pos: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast le_pos_low: {e}")))?;
|
||||
let flat = ge_neg_low
|
||||
.mul(&le_pos_low)
|
||||
.map_err(|e| MLError::ModelError(format!("mul flat: {e}")))?;
|
||||
|
||||
// Long50: pred > +low AND pred <= +high
|
||||
let gt_pos_low = pred
|
||||
.gt(&low_pos)
|
||||
.map_err(|e| MLError::ModelError(format!("gt low_pos: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast gt_pos_low: {e}")))?;
|
||||
let le_pos_high = pred
|
||||
.le(&high_pos)
|
||||
.map_err(|e| MLError::ModelError(format!("le high_pos: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast le_pos_high: {e}")))?;
|
||||
let long50 = gt_pos_low
|
||||
.mul(&le_pos_high)
|
||||
.map_err(|e| MLError::ModelError(format!("mul long50: {e}")))?;
|
||||
|
||||
// Long100: pred > +high
|
||||
let long100 = pred
|
||||
.gt(&high_pos)
|
||||
.map_err(|e| MLError::ModelError(format!("gt high_pos: {e}")))?
|
||||
.to_dtype(DType::F32)
|
||||
.map_err(|e| MLError::ModelError(format!("cast long100: {e}")))?;
|
||||
|
||||
// Stack [batch] tensors into [batch, 5]
|
||||
Tensor::stack(&[short100, short50, flat, long50, long100], 1)
|
||||
.map_err(|e| MLError::ModelError(format!("stack action scores: {e}")))
|
||||
Ok(out)
|
||||
}
|
||||
|
||||
/// Extract the median signal from TFT quantile predictions.
|
||||
///
|
||||
/// Input shape: `[batch, horizon, num_quantiles]` (typically `[batch, 1, 3]`).
|
||||
/// Returns `[batch]` — the median (quantile index 1) from the first horizon step.
|
||||
/// Input layout: `[batch * horizon * num_quantiles]` (typically `[batch * 1 * 3]`).
|
||||
/// Returns `[batch]` -- the median (quantile index 1) from the first horizon step.
|
||||
///
|
||||
/// # Arguments
|
||||
/// * `quantiles` - GPU-resident `[batch * horizon * num_quantiles]` data
|
||||
/// * `batch` - batch size
|
||||
/// * `horizon` - number of horizon steps (>= 1)
|
||||
/// * `num_quantiles` - number of quantiles (>= 2, median at index 1)
|
||||
/// * `stream` - CUDA stream for kernel launch and allocation
|
||||
///
|
||||
/// # Returns
|
||||
/// GPU-resident `CudaSlice<f32>` of shape `[batch]`.
|
||||
///
|
||||
/// # Errors
|
||||
/// Returns `MLError::ModelError` if the tensor has fewer than 3 dimensions
|
||||
/// or the quantile/horizon dimensions are too small.
|
||||
pub fn tft_quantile_to_signal(quantiles: &Tensor) -> Result<Tensor, MLError> {
|
||||
let shape = quantiles.dims();
|
||||
if shape.len() < 3 {
|
||||
return Err(MLError::ModelError(format!(
|
||||
"tft_quantile_to_signal: expected 3 dims, got {}",
|
||||
shape.len()
|
||||
)));
|
||||
}
|
||||
let horizon = shape.get(1).copied().unwrap_or(0);
|
||||
/// Returns `MLError::ModelError` if dimensions are invalid or kernel fails.
|
||||
pub fn tft_quantile_to_signal(
|
||||
quantiles: &CudaSlice<f32>,
|
||||
batch: usize,
|
||||
horizon: usize,
|
||||
num_quantiles: usize,
|
||||
stream: &CudaStream,
|
||||
) -> Result<CudaSlice<f32>, MLError> {
|
||||
if horizon < 1 {
|
||||
return Err(MLError::ModelError(
|
||||
"tft_quantile_to_signal: horizon dim must be >= 1".to_string(),
|
||||
"tft_quantile_to_signal: horizon must be >= 1".to_owned(),
|
||||
));
|
||||
}
|
||||
let num_q = shape.get(2).copied().unwrap_or(0);
|
||||
if num_q < 2 {
|
||||
if num_quantiles < 2 {
|
||||
return Err(MLError::ModelError(format!(
|
||||
"tft_quantile_to_signal: need >= 2 quantiles, got {num_q}"
|
||||
"tft_quantile_to_signal: need >= 2 quantiles, got {num_quantiles}"
|
||||
)));
|
||||
}
|
||||
|
||||
// narrow(dim=1, start=0, len=1) -> first horizon step
|
||||
// narrow(dim=2, start=1, len=1) -> median quantile
|
||||
// squeeze both singleton dims
|
||||
quantiles
|
||||
.narrow(1, 0, 1)
|
||||
.map_err(|e| MLError::ModelError(format!("narrow horizon: {e}")))?
|
||||
.narrow(2, 1, 1)
|
||||
.map_err(|e| MLError::ModelError(format!("narrow quantile: {e}")))?
|
||||
.squeeze(2)
|
||||
.map_err(|e| MLError::ModelError(format!("squeeze quantile: {e}")))?
|
||||
.squeeze(1)
|
||||
.map_err(|e| MLError::ModelError(format!("squeeze horizon: {e}")))
|
||||
}
|
||||
let context = stream.context();
|
||||
let kernels = load_kernels(&context)?;
|
||||
|
||||
/// Run GPU backtest for a supervised model using signal thresholds.
|
||||
///
|
||||
/// Common helper for all 8 supervised hyperopt adapters.
|
||||
/// Takes a forward function, validation data, and thresholds.
|
||||
/// Returns `(sharpe, total_trades, threshold_distance)` or error.
|
||||
pub fn evaluate_supervised_gpu_backtest(
|
||||
val_features: &[Vec<f32>],
|
||||
val_prices: &[[f32; 4]],
|
||||
feature_dim: usize,
|
||||
signal_high_bps: f32,
|
||||
signal_low_bps: f32,
|
||||
max_position: f32,
|
||||
device: &candle_core::Device,
|
||||
forward_fn: &dyn Fn(&Tensor) -> Result<Tensor, MLError>,
|
||||
) -> Result<(f32, u32, f32), MLError> {
|
||||
use super::gpu_backtest_evaluator::{GpuBacktestConfig, GpuBacktestEvaluator};
|
||||
let mut out = stream
|
||||
.alloc_zeros::<f32>(batch)
|
||||
.map_err(|e| MLError::ModelError(format!("alloc tft signal: {e}")))?;
|
||||
|
||||
if val_features.len() < 100 {
|
||||
return Err(MLError::ConfigError(format!(
|
||||
"evaluate_supervised_gpu_backtest: need >= 100 validation rows, got {}",
|
||||
val_features.len()
|
||||
)));
|
||||
}
|
||||
|
||||
let config = GpuBacktestConfig {
|
||||
max_position,
|
||||
tx_cost_bps: 2.0,
|
||||
spread_cost: 0.5,
|
||||
initial_capital: 100_000.0,
|
||||
..Default::default()
|
||||
let batch_i32 = batch as i32;
|
||||
let horizon_i32 = horizon as i32;
|
||||
let num_q_i32 = num_quantiles as i32;
|
||||
let grid = ((batch as u32 + 255) / 256).max(1);
|
||||
let cfg = LaunchConfig {
|
||||
grid_dim: (grid, 1, 1),
|
||||
block_dim: (256, 1, 1),
|
||||
shared_mem_bytes: 0,
|
||||
};
|
||||
|
||||
// Wrap data as single-window slices for GpuBacktestEvaluator
|
||||
let prices_window: Vec<[f32; 4]> = val_prices.to_vec();
|
||||
let features_window: Vec<Vec<f32>> = val_features.to_vec();
|
||||
// Safety: argument order matches tft_quantile_extract_kernel(
|
||||
// quantiles, out_signal, horizon, num_quantiles, batch)
|
||||
unsafe {
|
||||
stream
|
||||
.launch_builder(&kernels.tft_extract)
|
||||
.arg(quantiles)
|
||||
.arg(&mut out)
|
||||
.arg(&horizon_i32)
|
||||
.arg(&num_q_i32)
|
||||
.arg(&batch_i32)
|
||||
.launch(cfg)
|
||||
.map_err(|e| MLError::ModelError(format!("tft_quantile_extract launch: {e}")))?;
|
||||
}
|
||||
|
||||
let mut evaluator = GpuBacktestEvaluator::new(
|
||||
&[prices_window],
|
||||
&[features_window],
|
||||
feature_dim,
|
||||
config,
|
||||
device,
|
||||
)?;
|
||||
|
||||
let portfolio_dim = 3;
|
||||
|
||||
let metrics = evaluator.evaluate(
|
||||
&|states: &Tensor| {
|
||||
let pred = forward_fn(states)?;
|
||||
// Handle [batch, 1] vs [batch] output shape
|
||||
let signal = if pred.dims().len() == 2 {
|
||||
let last_dim = pred.dims().get(1).copied().unwrap_or(0);
|
||||
if last_dim == 1 {
|
||||
pred.squeeze(1)
|
||||
.map_err(|e| MLError::ModelError(format!("squeeze pred: {e}")))?
|
||||
} else {
|
||||
pred
|
||||
}
|
||||
} else {
|
||||
pred
|
||||
};
|
||||
// Convert scalar signals to [batch, 5] action scores via thresholds
|
||||
signal_to_action_scores(&signal, signal_high_bps, signal_low_bps)
|
||||
},
|
||||
portfolio_dim,
|
||||
device,
|
||||
)?;
|
||||
|
||||
let first = metrics.first().ok_or_else(|| {
|
||||
MLError::ModelError("evaluate_supervised_gpu_backtest: no window metrics returned".into())
|
||||
})?;
|
||||
|
||||
let sharpe = first.sharpe;
|
||||
let total_trades = first.total_trades as u32;
|
||||
|
||||
let threshold_distance =
|
||||
((signal_high_bps - 10.0).abs() / 10.0) + ((signal_low_bps - 5.0).abs() / 5.0);
|
||||
|
||||
Ok((sharpe, total_trades, threshold_distance))
|
||||
Ok(out)
|
||||
}
|
||||
|
||||
/// Compute walk-forward backtest fitness for hyperopt (minimization).
|
||||
@@ -300,123 +260,177 @@ pub fn backtest_fitness(
|
||||
mod tests {
|
||||
use super::*;
|
||||
|
||||
fn cuda_device() -> candle_core::Device {
|
||||
candle_core::Device::new_cuda(0).expect("CUDA device required")
|
||||
fn cuda_stream() -> CudaStream {
|
||||
let dev = candle_core::Device::new_cuda(0).expect("CUDA device required");
|
||||
match dev {
|
||||
candle_core::Device::Cuda(d) => d.cuda_stream().clone(),
|
||||
_ => panic!("expected CUDA device"),
|
||||
}
|
||||
}
|
||||
|
||||
// ── ppo_to_exposure_scores ──────────────────────────────────────────
|
||||
|
||||
#[test]
|
||||
fn test_ppo_to_exposure_scores_shape() {
|
||||
let stream = cuda_stream();
|
||||
let batch = 4;
|
||||
let uniform = vec![1.0_f32 / 45.0; batch * 45];
|
||||
let probs = Tensor::from_vec(uniform, &[batch, 45], &cuda_device()).unwrap();
|
||||
let scores = ppo_to_exposure_scores(&probs).unwrap();
|
||||
assert_eq!(scores.dims(), &[batch, 5]);
|
||||
let mut probs_buf = stream.alloc_zeros::<f32>(batch * 45).unwrap();
|
||||
stream.memcpy_htod(&uniform, &mut probs_buf).unwrap();
|
||||
|
||||
// GPU-side: all values should be ~0.2 (each of 5 bins sums 9 cells of 1/45)
|
||||
let expected = Tensor::new(0.2_f32, scores.device()).unwrap();
|
||||
let max_diff = scores.broadcast_sub(&expected).unwrap().abs().unwrap()
|
||||
.max(0).unwrap().max(0).unwrap().to_scalar::<f32>().unwrap();
|
||||
assert!(max_diff < 1e-5, "expected ~0.2 everywhere, max_diff={max_diff}");
|
||||
let scores = ppo_to_exposure_scores(&probs_buf, batch, &stream).unwrap();
|
||||
|
||||
// Download and check shape + values
|
||||
let mut host = vec![0.0_f32; batch * 5];
|
||||
stream.memcpy_dtoh(&scores, &mut host).unwrap();
|
||||
stream.synchronize().unwrap();
|
||||
|
||||
assert_eq!(host.len(), batch * 5);
|
||||
// Each of 5 bins sums 9 cells of 1/45 = 0.2
|
||||
for val in &host {
|
||||
assert!(
|
||||
(*val - 0.2).abs() < 1e-5,
|
||||
"expected ~0.2 everywhere, got {val}"
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_ppo_to_exposure_scores_argmax() {
|
||||
// Put all mass on actions 36..44 → exposure bin 4 (Long100)
|
||||
let stream = cuda_stream();
|
||||
// Put all mass on actions 36..44 => exposure bin 4 (Long100)
|
||||
let batch = 2;
|
||||
let mut raw = vec![0.0_f32; batch * 45];
|
||||
for b in 0..batch {
|
||||
// actions 36..44 = exposure bucket 4, each 1/9
|
||||
for a in 36..45 {
|
||||
if let Some(slot) = raw.get_mut(b * 45 + a) {
|
||||
*slot = 1.0 / 9.0;
|
||||
}
|
||||
}
|
||||
}
|
||||
let probs = Tensor::from_vec(raw, &[batch, 45], &cuda_device()).unwrap();
|
||||
let scores = ppo_to_exposure_scores(&probs).unwrap();
|
||||
// GPU-side argmax: all rows should have argmax at index 4 (Long100)
|
||||
let argmaxes = scores.argmax(1).unwrap();
|
||||
let expected_4 = Tensor::new(&[4_u32, 4_u32], scores.device()).unwrap();
|
||||
let all_match = argmaxes.eq(&expected_4).unwrap()
|
||||
.to_dtype(DType::U32).unwrap()
|
||||
.sum_all().unwrap().to_scalar::<u32>().unwrap();
|
||||
assert_eq!(all_match, batch as u32, "all rows should argmax to Long100 (4)");
|
||||
}
|
||||
let mut probs_buf = stream.alloc_zeros::<f32>(batch * 45).unwrap();
|
||||
stream.memcpy_htod(&raw, &mut probs_buf).unwrap();
|
||||
|
||||
#[test]
|
||||
fn test_ppo_to_exposure_scores_rejects_wrong_shape() {
|
||||
let probs = Tensor::zeros(&[8, 5], DType::F32, &cuda_device()).unwrap();
|
||||
let err = ppo_to_exposure_scores(&probs);
|
||||
assert!(err.is_err(), "should reject [8,5]");
|
||||
let scores = ppo_to_exposure_scores(&probs_buf, batch, &stream).unwrap();
|
||||
let mut host = vec![0.0_f32; batch * 5];
|
||||
stream.memcpy_dtoh(&scores, &mut host).unwrap();
|
||||
stream.synchronize().unwrap();
|
||||
|
||||
// Each row should have argmax at index 4 (Long100)
|
||||
for b in 0..batch {
|
||||
let row_start = b * 5;
|
||||
let mut best_idx = 0;
|
||||
let mut best_val = f32::NEG_INFINITY;
|
||||
for i in 0..5 {
|
||||
let v = host[row_start + i];
|
||||
if v > best_val {
|
||||
best_val = v;
|
||||
best_idx = i;
|
||||
}
|
||||
}
|
||||
assert_eq!(best_idx, 4, "batch {b}: expected argmax=4, got {best_idx}");
|
||||
}
|
||||
}
|
||||
|
||||
// ── signal_to_action_scores ─────────────────────────────────────────
|
||||
|
||||
/// GPU-side argmax for a single-row tensor.
|
||||
fn gpu_argmax(t: &Tensor) -> u32 {
|
||||
t.argmax(1).unwrap()
|
||||
.flatten_all().unwrap()
|
||||
.squeeze(0).unwrap()
|
||||
.to_scalar::<u32>().unwrap()
|
||||
/// Host-side argmax over a [1, 5] row downloaded from GPU.
|
||||
fn host_argmax(host: &[f32]) -> usize {
|
||||
let mut best_idx = 0;
|
||||
let mut best_val = f32::NEG_INFINITY;
|
||||
for (i, &v) in host.iter().enumerate() {
|
||||
if v > best_val {
|
||||
best_val = v;
|
||||
best_idx = i;
|
||||
}
|
||||
}
|
||||
best_idx
|
||||
}
|
||||
|
||||
fn run_signal_test(pred_val: f32, high: f32, low: f32) -> Vec<f32> {
|
||||
let stream = cuda_stream();
|
||||
let mut pred_buf = stream.alloc_zeros::<f32>(1).unwrap();
|
||||
stream.memcpy_htod(&[pred_val], &mut pred_buf).unwrap();
|
||||
let scores = signal_to_action_scores(&pred_buf, 1, high, low, &stream).unwrap();
|
||||
let mut host = vec![0.0_f32; 5];
|
||||
stream.memcpy_dtoh(&scores, &mut host).unwrap();
|
||||
stream.synchronize().unwrap();
|
||||
host
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_signal_to_action_scores_strong_long() {
|
||||
let pred = Tensor::new(&[20.0_f32], &cuda_device()).unwrap();
|
||||
let scores = signal_to_action_scores(&pred, 10.0, 5.0).unwrap();
|
||||
assert_eq!(scores.dims(), &[1, 5]);
|
||||
assert_eq!(gpu_argmax(&scores), 4); // Long100
|
||||
let host = run_signal_test(20.0, 10.0, 5.0);
|
||||
assert_eq!(host_argmax(&host), 4); // Long100
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_signal_to_action_scores_strong_short() {
|
||||
let pred = Tensor::new(&[-20.0_f32], &cuda_device()).unwrap();
|
||||
let scores = signal_to_action_scores(&pred, 10.0, 5.0).unwrap();
|
||||
assert_eq!(gpu_argmax(&scores), 0); // Short100
|
||||
let host = run_signal_test(-20.0, 10.0, 5.0);
|
||||
assert_eq!(host_argmax(&host), 0); // Short100
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_signal_to_action_scores_flat() {
|
||||
let pred = Tensor::new(&[0.0_f32], &cuda_device()).unwrap();
|
||||
let scores = signal_to_action_scores(&pred, 10.0, 5.0).unwrap();
|
||||
assert_eq!(gpu_argmax(&scores), 2); // Flat
|
||||
let host = run_signal_test(0.0, 10.0, 5.0);
|
||||
assert_eq!(host_argmax(&host), 2); // Flat
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_signal_to_action_scores_mild_long() {
|
||||
let pred = Tensor::new(&[7.0_f32], &cuda_device()).unwrap();
|
||||
let scores = signal_to_action_scores(&pred, 10.0, 5.0).unwrap();
|
||||
assert_eq!(gpu_argmax(&scores), 3); // Long50
|
||||
let host = run_signal_test(7.0, 10.0, 5.0);
|
||||
assert_eq!(host_argmax(&host), 3); // Long50
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_signal_to_action_scores_mild_short() {
|
||||
let pred = Tensor::new(&[-7.0_f32], &cuda_device()).unwrap();
|
||||
let scores = signal_to_action_scores(&pred, 10.0, 5.0).unwrap();
|
||||
assert_eq!(gpu_argmax(&scores), 1); // Short50
|
||||
let host = run_signal_test(-7.0, 10.0, 5.0);
|
||||
assert_eq!(host_argmax(&host), 1); // Short50
|
||||
}
|
||||
|
||||
// ── tft_quantile_to_signal ──────────────────────────────────────────
|
||||
|
||||
#[test]
|
||||
fn test_tft_quantile_to_signal() {
|
||||
let stream = cuda_stream();
|
||||
// [batch=2, horizon=1, quantiles=3]
|
||||
// quantiles: [q10, median, q90]
|
||||
let data: Vec<f32> = vec![
|
||||
-1.0, 0.5, 2.0, // batch 0: median = 0.5
|
||||
-3.0, -1.5, 0.0, // batch 1: median = -1.5
|
||||
-1.0, 0.5, 2.0, // batch 0: median = 0.5
|
||||
-3.0, -1.5, 0.0, // batch 1: median = -1.5
|
||||
];
|
||||
let q = Tensor::from_vec(data, &[2, 1, 3], &cuda_device()).unwrap();
|
||||
let signal = tft_quantile_to_signal(&q).unwrap();
|
||||
assert_eq!(signal.dims(), &[2]);
|
||||
let mut q_buf = stream.alloc_zeros::<f32>(data.len()).unwrap();
|
||||
stream.memcpy_htod(&data, &mut q_buf).unwrap();
|
||||
|
||||
// GPU-side comparison: signal should equal [0.5, -1.5]
|
||||
let expected = Tensor::new(&[0.5_f32, -1.5], signal.device()).unwrap();
|
||||
let max_diff = signal.sub(&expected).unwrap().abs().unwrap()
|
||||
.max(0).unwrap().to_scalar::<f32>().unwrap();
|
||||
assert!(max_diff < 1e-6, "tft signal mismatch: max_diff={max_diff}");
|
||||
let signal = tft_quantile_to_signal(&q_buf, 2, 1, 3, &stream).unwrap();
|
||||
let mut host = vec![0.0_f32; 2];
|
||||
stream.memcpy_dtoh(&signal, &mut host).unwrap();
|
||||
stream.synchronize().unwrap();
|
||||
|
||||
assert!(
|
||||
(host[0] - 0.5).abs() < 1e-6,
|
||||
"batch 0: expected 0.5, got {}",
|
||||
host[0]
|
||||
);
|
||||
assert!(
|
||||
(host[1] - (-1.5)).abs() < 1e-6,
|
||||
"batch 1: expected -1.5, got {}",
|
||||
host[1]
|
||||
);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn test_tft_quantile_rejects_bad_dims() {
|
||||
let stream = cuda_stream();
|
||||
let mut q_buf = stream.alloc_zeros::<f32>(6).unwrap();
|
||||
stream.memcpy_htod(&[0.0_f32; 6], &mut q_buf).unwrap();
|
||||
|
||||
// horizon=0 should fail
|
||||
let err = tft_quantile_to_signal(&q_buf, 2, 0, 3, &stream);
|
||||
assert!(err.is_err(), "should reject horizon=0");
|
||||
|
||||
// num_quantiles=1 should fail
|
||||
let err = tft_quantile_to_signal(&q_buf, 2, 1, 1, &stream);
|
||||
assert!(err.is_err(), "should reject num_quantiles=1");
|
||||
}
|
||||
|
||||
// ── backtest_fitness ────────────────────────────────────────────────
|
||||
|
||||
95
crates/ml/src/cuda_pipeline/signal_adapter_kernel.cu
Normal file
95
crates/ml/src/cuda_pipeline/signal_adapter_kernel.cu
Normal file
@@ -0,0 +1,95 @@
|
||||
/**
|
||||
* Signal adapter CUDA kernels for converting model outputs to trading signals.
|
||||
*
|
||||
* Three fused kernels replacing Candle tensor operations:
|
||||
*
|
||||
* 1. ppo_to_exposure_scores_kernel:
|
||||
* Aggregates PPO 45-action softmax probs into 5 exposure scores.
|
||||
* Input: [batch * 45] float probs
|
||||
* Output: [batch * 5] float scores (sum of 9 order/urgency combos per bucket)
|
||||
*
|
||||
* 2. signal_to_action_scores_kernel:
|
||||
* Maps scalar predictions to [batch, 5] one-hot-ish action scores via thresholds.
|
||||
* Input: [batch] float predictions
|
||||
* Output: [batch * 5] float scores (one-hot based on threshold buckets)
|
||||
*
|
||||
* 3. tft_quantile_extract_kernel:
|
||||
* Extracts median signal from TFT quantile predictions.
|
||||
* Input: [batch * horizon * num_quantiles] float
|
||||
* Output: [batch] float median values (quantile index 1, horizon index 0)
|
||||
*
|
||||
* None of these kernels require common_device_functions.cuh (standalone).
|
||||
* Launch config: grid=(ceil(batch/256), 1, 1), block=(256, 1, 1).
|
||||
* One thread per batch element.
|
||||
*/
|
||||
|
||||
extern "C" __global__ void ppo_to_exposure_scores_kernel(
|
||||
const float* __restrict__ probs, /* [batch * 45] softmax probs */
|
||||
float* __restrict__ out_scores, /* [batch * 5] exposure scores */
|
||||
int batch
|
||||
) {
|
||||
int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= batch) return;
|
||||
|
||||
/* For each of 5 exposure buckets, sum over 9 order/urgency combos.
|
||||
* probs layout: [batch, 45] where 45 = 5 exposure * 9 (3 order * 3 urgency)
|
||||
* bucket e covers indices [e*9 .. e*9+8] within each batch row. */
|
||||
int base = b * 45;
|
||||
for (int e = 0; e < 5; ++e) {
|
||||
float sum = 0.0f;
|
||||
int offset = base + e * 9;
|
||||
for (int j = 0; j < 9; ++j) {
|
||||
sum += probs[offset + j];
|
||||
}
|
||||
out_scores[b * 5 + e] = sum;
|
||||
}
|
||||
}
|
||||
|
||||
extern "C" __global__ void signal_to_action_scores_kernel(
|
||||
const float* __restrict__ predictions, /* [batch] scalar signals */
|
||||
float* __restrict__ out_scores, /* [batch * 5] one-hot scores */
|
||||
float high_threshold_bps,
|
||||
float low_threshold_bps,
|
||||
int batch
|
||||
) {
|
||||
int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= batch) return;
|
||||
|
||||
float pred = predictions[b];
|
||||
int out_base = b * 5;
|
||||
|
||||
/* Zero all 5 scores, then set the matching bucket to 1.0 */
|
||||
out_scores[out_base + 0] = 0.0f;
|
||||
out_scores[out_base + 1] = 0.0f;
|
||||
out_scores[out_base + 2] = 0.0f;
|
||||
out_scores[out_base + 3] = 0.0f;
|
||||
out_scores[out_base + 4] = 0.0f;
|
||||
|
||||
if (pred < -high_threshold_bps) {
|
||||
out_scores[out_base + 0] = 1.0f; /* Short100 */
|
||||
} else if (pred < -low_threshold_bps) {
|
||||
out_scores[out_base + 1] = 1.0f; /* Short50 */
|
||||
} else if (pred > high_threshold_bps) {
|
||||
out_scores[out_base + 4] = 1.0f; /* Long100 */
|
||||
} else if (pred > low_threshold_bps) {
|
||||
out_scores[out_base + 3] = 1.0f; /* Long50 */
|
||||
} else {
|
||||
out_scores[out_base + 2] = 1.0f; /* Flat */
|
||||
}
|
||||
}
|
||||
|
||||
extern "C" __global__ void tft_quantile_extract_kernel(
|
||||
const float* __restrict__ quantiles, /* [batch * horizon * num_q] */
|
||||
float* __restrict__ out_signal, /* [batch] median values */
|
||||
int horizon,
|
||||
int num_quantiles,
|
||||
int batch
|
||||
) {
|
||||
int b = blockIdx.x * blockDim.x + threadIdx.x;
|
||||
if (b >= batch) return;
|
||||
|
||||
/* Extract quantile index 1 (median) from horizon step 0.
|
||||
* Layout: quantiles[b * horizon * num_q + 0 * num_q + 1] */
|
||||
int idx = b * horizon * num_quantiles + 1; /* horizon=0, quantile=1 */
|
||||
out_signal[b] = quantiles[idx];
|
||||
}
|
||||
@@ -52,7 +52,9 @@ use crate::ppo::trajectory_replay::TrajectoryReplayBuffer;
|
||||
use crate::MLError;
|
||||
|
||||
use crate::cuda_pipeline::gpu_backtest_evaluator::{GpuBacktestConfig, GpuBacktestEvaluator};
|
||||
use crate::cuda_pipeline::gpu_action_selector::cuda_f32_to_tensor;
|
||||
use crate::cuda_pipeline::signal_adapter::ppo_to_exposure_scores;
|
||||
use candle_core::cuda_backend::cudarc::driver::CudaSlice;
|
||||
|
||||
/// Pure model VRAM in MB (actor + critic + optimizers + gradients).
|
||||
const MODEL_OVERHEAD_MB: f64 = 300.0;
|
||||
@@ -1324,11 +1326,26 @@ impl PPOTrainer {
|
||||
device,
|
||||
)?;
|
||||
|
||||
// Forward function: PPO actor → softmax → ppo_to_exposure_scores → [B, 5]
|
||||
// Forward function: PPO actor -> softmax -> ppo_to_exposure_scores -> [B, 5]
|
||||
// Bridge: extract CudaSlice -> pure cudarc kernel -> DtoD copy back to Tensor.
|
||||
let metrics = evaluator.evaluate(
|
||||
&|states: &candle_core::Tensor| -> Result<candle_core::Tensor, MLError> {
|
||||
let probs = ppo.actor.action_probabilities(states)?;
|
||||
ppo_to_exposure_scores(&probs)
|
||||
let batch = probs.dims().first().copied().unwrap_or(1);
|
||||
let (storage_guard, _layout) = probs.storage_and_layout();
|
||||
let cuda_slice: &CudaSlice<f32> = match &*storage_guard {
|
||||
candle_core::Storage::Cuda(cs) => cs.as_cuda_slice()
|
||||
.map_err(|e| MLError::ModelError(format!("probs as_cuda_slice: {e}")))?,
|
||||
_ => return Err(MLError::ModelError("probs not on CUDA".into())),
|
||||
};
|
||||
let cuda_dev = match device {
|
||||
Device::Cuda(ref d) => d,
|
||||
_ => return Err(MLError::ModelError("device is not CUDA".into())),
|
||||
};
|
||||
let stream = cuda_dev.cuda_stream();
|
||||
let scores_slice = ppo_to_exposure_scores(cuda_slice, batch, &stream)?;
|
||||
drop(storage_guard);
|
||||
cuda_f32_to_tensor(&scores_slice, &[batch, 5], device)
|
||||
},
|
||||
3, // portfolio_dim
|
||||
device,
|
||||
|
||||
Reference in New Issue
Block a user