diff --git a/crates/ml-core/src/cuda_autograd/reduction_kernels.cu b/crates/ml-core/src/cuda_autograd/reduction_kernels.cu
index 36fec171e..4e160dfa1 100644
--- a/crates/ml-core/src/cuda_autograd/reduction_kernels.cu
+++ b/crates/ml-core/src/cuda_autograd/reduction_kernels.cu
@@ -87,37 +87,37 @@ void fused_stats_reduce(const float* __restrict__ data,
         w_cnt = warp_reduce_sum_bf16(w_cnt);
 
         if (threadIdx.x == 0) {
-            // atomicMin for bf16 via CAS loop
+            // atomicMin for f32 via CAS loop (unsigned int = 4 bytes = sizeof(float))
             {
-                unsigned short* addr_us = (unsigned short*)&result[0];
-                unsigned short old_bits = *addr_us;
-                unsigned short new_bits = *(unsigned short*)&w_min;
-                while (*(float*)&old_bits > w_min) {
-                    unsigned short assumed = old_bits;
-                    old_bits = atomicCAS(addr_us, assumed, new_bits);
+                unsigned int* addr_ui = (unsigned int*)&result[0];
+                unsigned int old_bits = *addr_ui;
+                unsigned int new_bits = __float_as_uint(w_min);
+                while (__uint_as_float(old_bits) > w_min) {
+                    unsigned int assumed = old_bits;
+                    old_bits = atomicCAS(addr_ui, assumed, new_bits);
                     if (old_bits == assumed) break;
-                    float cur = *(float*)&old_bits;
-                    float new_val = bf16_fmin(cur, w_min);
-                    new_bits = *(unsigned short*)&new_val;
+                    float cur = __uint_as_float(old_bits);
+                    float new_val = fminf(cur, w_min);
+                    new_bits = __float_as_uint(new_val);
                 }
             }
-            // atomicMax for bf16 via CAS loop
+            // atomicMax for f32 via CAS loop
             {
-                unsigned short* addr_us = (unsigned short*)&result[1];
-                unsigned short old_bits = *addr_us;
-                unsigned short new_bits = *(unsigned short*)&w_max;
-                while (*(float*)&old_bits < w_max) {
-                    unsigned short assumed = old_bits;
-                    old_bits = atomicCAS(addr_us, assumed, new_bits);
+                unsigned int* addr_ui = (unsigned int*)&result[1];
+                unsigned int old_bits = *addr_ui;
+                unsigned int new_bits = __float_as_uint(w_max);
+                while (__uint_as_float(old_bits) < w_max) {
+                    unsigned int assumed = old_bits;
+                    old_bits = atomicCAS(addr_ui, assumed, new_bits);
                     if (old_bits == assumed) break;
-                    float cur = *(float*)&old_bits;
-                    float new_val = bf16_fmax(cur, w_max);
-                    new_bits = *(unsigned short*)&new_val;
+                    float cur = __uint_as_float(old_bits);
+                    float new_val = fmaxf(cur, w_max);
+                    new_bits = __float_as_uint(new_val);
                 }
             }
-            atomicAddBF16(&result[2], w_sum);
-            atomicAddBF16(&result[3], w_ssq);
-            atomicAddBF16(&result[4], w_cnt);
+            atomicAdd(&result[2], w_sum);
+            atomicAdd(&result[3], w_ssq);
+            atomicAdd(&result[4], w_cnt);
         }
     }
 }
@@ -187,14 +187,14 @@ void argmax_flat(const float* __restrict__ data,
 
         if (threadIdx.x == 0) {
             // CAS loop to merge block results into global output
-            // result[0] = best val (as bf16), result[1] = best idx (as bf16)
-            unsigned short* val_addr = (unsigned short*)&result[0];
-            unsigned short old_bits = *val_addr;
-            while (*(float*)&old_bits < w_val) {
-                unsigned short assumed = old_bits;
-                old_bits = atomicCAS(val_addr, assumed, *(unsigned short*)&w_val);
+            // result[0] = best val (as f32), result[1] = best idx (as f32)
+            unsigned int* val_addr = (unsigned int*)&result[0];
+            unsigned int old_bits = *val_addr;
+            while (__uint_as_float(old_bits) < w_val) {
+                unsigned int assumed = old_bits;
+                old_bits = atomicCAS(val_addr, assumed, __float_as_uint(w_val));
                 if (old_bits == assumed) {
-                    result[1] = bf16((int)w_idx);
+                    result[1] = (float)w_idx;
                     break;
                 }
             }
diff --git a/crates/ml-core/src/cuda_autograd/reductions.rs b/crates/ml-core/src/cuda_autograd/reductions.rs
index b95e24152..d101c1990 100644
--- a/crates/ml-core/src/cuda_autograd/reductions.rs
+++ b/crates/ml-core/src/cuda_autograd/reductions.rs
@@ -94,7 +94,7 @@ impl ReductionKernels {
         // Cap blocks to avoid excessive atomic contention
         let blocks = blocks.min(1024);
         let n_i32 = n as i32;
-        let shared_bytes = threads * 5 * 2; // 5 f32 per thread for shared reduction (min, max, sum, sum_sq, count)
+        let shared_bytes = threads * 5 * 4; // 5 f32 per thread for shared reduction (min, max, sum, sum_sq, count)
 
         let cfg = LaunchConfig {
             grid_dim: (blocks, 1, 1),
@@ -167,8 +167,8 @@ impl ReductionKernels {
         let blocks = ((n as u32) + threads - 1) / threads;
         let blocks = blocks.min(1024);
         let n_i32 = n as i32;
-        // Shared: [B f32 values (B*2 bytes)][B int indices (B*4 bytes)]
-        let shared_bytes = threads * 2 + threads * 4;
+        // Shared: [B f32 values (B*4 bytes)][B int indices (B*4 bytes)]
+        let shared_bytes = threads * 4 + threads * 4;
 
         let cfg = LaunchConfig {
             grid_dim: (blocks, 1, 1),
@@ -216,8 +216,8 @@ impl ReductionKernels {
         let threads = 256_u32.min(cols as u32).max(64);
         let rows_i32 = rows as i32;
         let cols_i32 = cols as i32;
-        // Shared: [B f32 values (B*2 bytes)][B int indices (B*4 bytes)]
-        let shared_bytes = threads * 2 + threads * 4;
+        // Shared: [B f32 values (B*4 bytes)][B int indices (B*4 bytes)]
+        let shared_bytes = threads * 4 + threads * 4;
 
         let cfg = LaunchConfig {
             grid_dim: (rows as u32, 1, 1),
@@ -267,7 +267,7 @@ impl ReductionKernels {
         let blocks = ((n as u32) + threads - 1) / threads;
         let blocks = blocks.min(1024);
         let n_i32 = n as i32;
-        let shared_bytes = threads * 2; // 1 f32 per thread
+        let shared_bytes = threads * 4; // 1 f32 (4 bytes) per thread
 
         let cfg = LaunchConfig {
             grid_dim: (blocks, 1, 1),
@@ -315,7 +315,7 @@ impl ReductionKernels {
         let threads = 256_u32.min(rows as u32);
         let rows_i32 = rows as i32;
         let cols_i32 = cols as i32;
-        let shared_bytes = threads * 2; // 1 f32 per thread
+        let shared_bytes = threads * 4; // 1 f32 (4 bytes) per thread
 
         let cfg = LaunchConfig {
             grid_dim: (cols as u32, 1, 1),
diff --git a/crates/ml-dqn/src/evaluation/engine.rs b/crates/ml-dqn/src/evaluation/engine.rs
index c18116357..27c625e2b 100644
--- a/crates/ml-dqn/src/evaluation/engine.rs
+++ b/crates/ml-dqn/src/evaluation/engine.rs
@@ -448,18 +448,19 @@ mod tests {
             engine.trades.first().map_or(false, |t| t.pnl > 0.0),
             "Price went up on long = profit"
         );
-        assert!((engine.current_exposure - 0.5).abs() < 1e-6);
+        assert!((engine.current_exposure - 0.25).abs() < 1e-6, "Long50 = Long×Small = 0.25 exposure");
     }
 
     #[test]
     fn factored_reversal_generates_trade() {
         let mut engine = EvaluationEngine::new(10000.0);
         engine.process_bar_factored(0, &bar(100.0), &market_action(ExposureLevel::Long100));
-        engine.process_bar_factored(1, &bar(105.0), &market_action(ExposureLevel::Short100));
+        engine.process_bar_factored(1, &bar(105.0), &market_action(ExposureLevel::Short50));
         assert!(
             !engine.trades.is_empty(),
             "Reversal should generate at least 1 trade"
         );
+        // Short50 = Short×Full = -1.0 exposure
         assert!((engine.current_exposure - (-1.0)).abs() < 1e-6);
     }