# BF16 Elimination — Pure f32/TF32 Training Pipeline

> **For agentic workers:** REQUIRED SUB-SKILL: Use superpowers:subagent-driven-development (recommended) or superpowers:executing-plans to implement this plan task-by-task. Steps use checkbox (`- [ ]`) syntax for tracking.

**Goal:** Eliminate ALL bf16 (2526 references, 54 cuda_pipeline files + 16 outer files) from the DQN training pipeline, establishing pure f32/TF32 everywhere. bf16's 8-bit mantissa destroys gradient precision at batch=16384, produces identical quantized grad norms, and is pointless when cublasLtMatmul already does f32→TF32 internally.

**Architecture:** Layered regex sweep + compiler-driven fix. Phase 1 rewrites the shared CUDA header and all .cu kernels (`float` replaces `__nv_bfloat16`). Phase 2 rewrites the Rust shared infrastructure (`mod.rs` helpers, `gpu_weights.rs` structs). Phase 3 rewrites the DQN trainer core. Phase 4 rewrites all consumers (experience collector, fused training, action selector, etc.). Phase 5 rewrites non-DQN modules (PPO, IQL, curiosity, decision transformer). Phase 6 cleans dead code and removes the `half` crate. Each phase compiles incrementally — regex does 80% of the work, the compiler catches the remaining 20%.

**Tech Stack:** Rust 1.85, CUDA 12.4/12.9, cudarc (vendored), cublasLtMatmul (f32 + TF32 tensor cores)

**Scope exclusion:** `fxcache.rs` bf16 is a DISK FORMAT for compact feature cache files — not GPU precision. It stays as-is.

---

## File Structure

| Phase | Files | Ref Count | Description |
|-------|-------|-----------|-------------|
| 1 | `common_device_functions.cuh` + 30 `.cu` kernels | ~1034 | CUDA bf16→float types |
| 2 | `mod.rs`, `gpu_weights.rs` | ~305 | Rust bf16 helpers + weight structs |
| 3 | `gpu_dqn_trainer.rs` | ~307 | DQN trainer struct, constructor, methods |
| 4 | 10 GPU modules (experience, fused, action, etc.) | ~400 | All DQN consumer modules |
| 5 | 8 non-DQN modules (PPO, IQL, DT, curiosity, etc.) | ~250 | Non-DQN trainers |
| 6 | 16 outer files + Cargo.toml | ~130 | Tests, config, dead code cleanup |

---

### Task 1: Rewrite `common_device_functions.cuh` — Kill bf16 Header

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/common_device_functions.cuh:1-114`

This is the foundation — every CUDA kernel includes this file via NVRTC concatenation. All `bf16_*` wrapper functions become identity passthrough (`float` in, `float` out) or direct `float` math calls.

- [ ] **Step 1: Replace the bf16 header section (lines 1-114)**

Replace the entire bf16 section with f32 equivalents. The key transformation:
- Delete `#include <cuda_bf16.h>`
- All `__nv_bfloat16` → `float`
- `bf16_sqrt(x)` → `sqrtf(x)` (inline identity)
- `bf16_zero()` → `0.0f`
- `bf16_one()` → `1.0f`
- `bf16(x)` → `(x)` (identity cast)
- `atomicAddBF16` → `atomicAdd` (native f32 atomicAdd exists on all SM)
- `bf16_shfl_xor/down` → `__shfl_xor_sync/__shfl_down_sync` (native f32)
- `bf16_warp_sum/max` → f32 warp reductions using `__shfl_xor_sync`
- `leaky_relu_bf16(x)` → inline f32 leaky relu
- `f32_to_bf16(x)` → `(x)` (identity)

```cuda
/* ------------------------------------------------------------------ */
/*  F32/TF32 Precision — All GPU buffers are native float.            */
/* ------------------------------------------------------------------ */
/* TF32 tensor cores (19-bit mantissa) activated via cublasLtMatmul   */
/* with CUBLAS_COMPUTE_32F. Storage is pure f32 everywhere.           */

/* ── F32 math wrappers (kept as thin aliases for kernel readability) ── */
__device__ __forceinline__ float bf16_sqrt(float x) { return sqrtf(x); }
__device__ __forceinline__ float bf16_log(float x) { return logf(x); }
__device__ __forceinline__ float bf16_exp(float x) { return expf(x); }
__device__ __forceinline__ float bf16_pow(float x, float p) { return powf(x, p); }
__device__ __forceinline__ float bf16_fabs(float x) { return fabsf(x); }
__device__ __forceinline__ float bf16_fmax(float a, float b) { return fmaxf(a, b); }
__device__ __forceinline__ float bf16_fmin(float a, float b) { return fminf(a, b); }
__device__ __forceinline__ float bf16_cos(float x) { return cosf(x); }
__device__ __forceinline__ float bf16_sin(float x) { return sinf(x); }
__device__ __forceinline__ float bf16_tanh(float x) { return tanhf(x); }
__device__ __forceinline__ float bf16_floor(float x) { return floorf(x); }
__device__ __forceinline__ float bf16_powf(float x, float p) { return powf(x, p); }

__device__ __forceinline__ float bf16_zero() { return 0.0f; }
__device__ __forceinline__ float bf16_one() { return 1.0f; }
__device__ __forceinline__ float bf16(float x) { return x; }

__device__ __forceinline__ float bf16_shfl_xor(unsigned mask, float val, int offset) {
    return __shfl_xor_sync(mask, val, offset);
}
__device__ __forceinline__ float bf16_shfl_down(unsigned mask, float val, int offset) {
    return __shfl_down_sync(mask, val, offset);
}
__device__ __forceinline__ float bf16_warp_sum(float val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val += __shfl_xor_sync(0xFFFFFFFF, val, offset);
    return val;
}
__device__ __forceinline__ float bf16_warp_max(float val) {
    for (int offset = 16; offset > 0; offset >>= 1)
        val = fmaxf(val, __shfl_xor_sync(0xFFFFFFFF, val, offset));
    return val;
}

__device__ __forceinline__ float leaky_relu_bf16(float x) {
    return (x > 0.0f) ? x : 0.01f * x;
}

__device__ __forceinline__ float f32_to_bf16(float x) { return x; }

/* atomicAdd for float is native on SM30+ — no CAS loop needed */
__device__ __forceinline__ void atomicAddBF16(float* addr, float val) {
    atomicAdd(addr, val);
}
```

- [ ] **Step 2: Compile check — verify header parses**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | head -30`
Expected: CUDA kernels that use `__nv_bfloat16` directly (not through wrappers) will fail. That's expected — Task 2 fixes those.

- [ ] **Step 3: Commit**

```bash
git add crates/ml/src/cuda_pipeline/common_device_functions.cuh
git commit -m "refactor: rewrite common_device_functions.cuh — bf16 wrappers now pure f32 identity functions"
```

---

### Task 2: Regex Sweep All CUDA Kernels — `__nv_bfloat16` → `float`

**Files:**
- Modify: All 30 `.cu` files in `crates/ml/src/cuda_pipeline/`

Every `.cu` file that uses `__nv_bfloat16` in function signatures, local variables, or shared memory needs mechanical replacement. The `common_device_functions.cuh` wrappers (bf16_zero, bf16_sqrt, etc.) already return `float` after Task 1, but kernel code that declares `__nv_bfloat16` variables directly must be converted.

- [ ] **Step 1: Run regex replacements on all .cu files**

Apply these sed patterns to every `.cu` file in `crates/ml/src/cuda_pipeline/`:

```bash
cd crates/ml/src/cuda_pipeline

# 1. Type declarations: __nv_bfloat16 → float
sed -i 's/__nv_bfloat16\b/float/g' *.cu

# 2. Explicit cast functions → identity (may leave redundant parens, that's fine)
sed -i 's/__float2bfloat16(\([^)]*\))/(\1)/g' *.cu
sed -i 's/__bfloat162float(\([^)]*\))/(\1)/g' *.cu

# 3. Remove cuda_bf16.h includes (already removed from .cuh, but some .cu may have it)
sed -i '/#include.*cuda_bf16\.h/d' *.cu
```

- [ ] **Step 2: Manual review of shared memory declarations**

Search for `__shared__ float` (was `__shared__ __nv_bfloat16`) — verify sizes are correct. Since f32 is 4 bytes vs bf16's 2 bytes, shared memory arrays might need size adjustments if they were sized in bytes. Most kernels size shared memory by element count, so this should be fine — but verify:

```bash
grep -n '__shared__.*float' *.cu | head -30
```

If any kernel used byte-sized shared memory math (e.g., `sizeof(__nv_bfloat16)`), fix the size to `sizeof(float)`.

- [ ] **Step 3: Fix atomicAddBF16 call sites**

After the regex, `atomicAddBF16` now takes `float*` (from the rewritten header). But call sites that cast pointers need fixing:

```bash
grep -n 'atomicAddBF16' *.cu
```

Change any `(float*)addr` patterns that were previously `(__nv_bfloat16*)addr`. Since everything is now float, the casts become unnecessary — remove them.

- [ ] **Step 4: Fix f32_to_bf16_kernel and bf16_to_f32_kernel in dqn_utility_kernels.cu**

These conversion kernels are DEAD CODE now (no bf16 buffers to convert). Either:
- Delete them entirely, OR
- Convert to identity kernels (f32→f32 copy) as placeholders until Rust references are cleaned up in Task 5

Recommended: Convert to identity (memcpy) kernels so Rust code compiles. They'll be deleted in Task 9.

```cuda
extern "C" __global__ void f32_to_bf16_kernel(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) dst[i] = src[i];
}

extern "C" __global__ void bf16_to_f32_kernel(
    const float* __restrict__ src,
    float* __restrict__ dst,
    int n)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) dst[i] = src[i];
}
```

Also convert `dqn_nan_check_bf16` to read `float*`:
```cuda
extern "C" __global__ void dqn_nan_check_bf16(
    const float* __restrict__ data,
    int n,
    int* __restrict__ flag)
```

- [ ] **Step 5: Compile check**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | head -50`
Expected: CUDA files compile clean (NVRTC runtime compilation), Rust errors will come from type mismatches (CudaSlice<half::bf16> vs f32 kernel args). That's expected — Tasks 3-6 fix the Rust side.

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/cuda_pipeline/*.cu
git commit -m "refactor: all CUDA kernels pure f32 — __nv_bfloat16 eliminated from 30 kernel files"
```

---

### Task 3: Rewrite `mod.rs` — Kill bf16 Transfer Helpers

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/mod.rs`

The shared bf16 helpers (`htod_f32_to_bf16`, `dtoh_bf16_to_f32`, `clone_htod_f32_to_bf16`, `f32_slice_to_bf16`, `buf_bytes_bf16`) are used across every GPU module. Replace them with direct f32 equivalents.

- [ ] **Step 1: Replace the 5 bf16 helper functions (lines 53-119)**

Replace the entire bf16 helper section with f32 direct transfer functions. Key transformations:

```rust
/// Byte size of an f32 buffer.
#[inline]
pub fn buf_bytes_f32(num_elements: usize) -> usize {
    num_elements * std::mem::size_of::<f32>()
}

/// Upload f32 host data into an existing `CudaSlice<f32>`.
pub fn htod_f32(
    stream: &Arc<CudaStream>,
    src: &[f32],
    dst: &mut CudaSlice<f32>,
) -> Result<(), MLError> {
    stream.memcpy_htod(src, dst)
        .map_err(|e| MLError::ModelError(format!("htod_f32: {e}")))
}

/// Upload f32 host data to a new `CudaSlice<f32>`.
pub fn clone_htod_f32(
    stream: &Arc<CudaStream>,
    src: &[f32],
) -> Result<CudaSlice<f32>, MLError> {
    stream.clone_htod(src)
        .map_err(|e| MLError::ModelError(format!("clone_htod_f32: {e}")))
}

/// Download `CudaSlice<f32>` to host f32 buffer.
pub fn dtoh_f32<Src: cudarc::driver::DevicePtr<f32>>(
    stream: &Arc<CudaStream>,
    src: &Src,
    dst: &mut [f32],
) -> Result<(), MLError> {
    let n = dst.len();
    let num_bytes = n * std::mem::size_of::<f32>();
    let (src_ptr, _guard) = src.device_ptr(stream);
    #[allow(unsafe_code)]
    unsafe {
        cudarc::driver::sys::cuStreamSynchronize(stream.cu_stream());
        cudarc::driver::sys::cuMemcpyDtoH_v2(
            dst.as_mut_ptr().cast(),
            src_ptr,
            num_bytes,
        );
    }
    Ok(())
}
```

Keep the old function names as `#[inline]` aliases that call the new functions, so all 188 call sites compile immediately:

```rust
/// DEPRECATED: Legacy alias — will be removed in cleanup phase.
#[inline]
pub fn htod_f32_to_bf16(
    stream: &Arc<CudaStream>,
    src: &[f32],
    dst: &mut CudaSlice<f32>,
) -> Result<(), MLError> {
    htod_f32(stream, src, dst)
}

#[inline]
pub fn clone_htod_f32_to_bf16(
    stream: &Arc<CudaStream>,
    src: &[f32],
) -> Result<CudaSlice<f32>, MLError> {
    clone_htod_f32(stream, src)
}

#[inline]
pub fn dtoh_bf16_to_f32<Src: cudarc::driver::DevicePtr<f32>>(
    stream: &Arc<CudaStream>,
    src: &Src,
    dst: &mut [f32],
) -> Result<(), MLError> {
    dtoh_f32(stream, src, dst)
}

#[inline]
pub fn f32_slice_to_bf16(src: &[f32]) -> Vec<f32> {
    src.to_vec()
}

#[inline]
pub fn buf_bytes_bf16(num_elements: usize) -> usize {
    buf_bytes_f32(num_elements)
}
```

- [ ] **Step 2: Convert GpuBatch struct fields (lines ~226-230)**

```rust
pub struct GpuBatch {
    pub features: CudaSlice<f32>,
    pub targets: CudaSlice<f32>,
    pub ofi_features: Option<CudaSlice<f32>>,
    // ... other fields stay the same
}
```

Update all GpuBatch constructors (`new`, `from_raw_gpu_data`) to use `clone_htod_f32` instead of `clone_htod_f32_to_bf16`.

- [ ] **Step 3: Convert `clone_device_buffer` function (~line 174)**

```rust
pub fn clone_device_buffer(
    stream: &Arc<CudaStream>,
    src: &CudaSlice<f32>,
) -> Result<CudaSlice<f32>, crate::MLError> {
    let n = src.len();
    let mut dst = stream.alloc_zeros::<f32>(n).map_err(|e| {
        crate::MLError::ModelError(format!("clone_device_buffer alloc({n}): {e}"))
    })?;
    unsafe {
        let num_bytes = n * std::mem::size_of::<f32>();
        cudarc::driver::sys::cuMemcpyDtoDAsync_v2(
            dst.raw_ptr(), src.raw_ptr(), num_bytes, stream.cu_stream()
        );
    }
    Ok(dst)
}
```

- [ ] **Step 4: Remove `use half;` and `half::bf16` imports**

```bash
grep -n 'use half' crates/ml/src/cuda_pipeline/mod.rs
```

Remove all half crate imports from mod.rs. Keep the `pub use` re-export ONLY if other modules still need it (they will until Tasks 4-8 are done — leave it for now).

- [ ] **Step 5: Compile check**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -30`
Expected: Errors in files that pass `CudaSlice<half::bf16>` to the now-f32 helper functions. This is the compiler telling us exactly which files need updating.

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/cuda_pipeline/mod.rs
git commit -m "refactor: mod.rs bf16 helpers → direct f32 transfers, legacy aliases kept for incremental migration"
```

---

### Task 4: Rewrite `gpu_weights.rs` — Kill bf16 Weight Structs

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_weights.rs`

All weight structs (`DqnWeightSet`, `CuriosityWeightSet`, `PpoWeightSet`, `CuriosityWeightSetBf16`) use `CudaSlice<half::bf16>`. Convert to `CudaSlice<f32>`.

- [ ] **Step 1: Regex replace all CudaSlice<half::bf16> → CudaSlice<f32>**

```bash
sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' crates/ml/src/cuda_pipeline/gpu_weights.rs
```

- [ ] **Step 2: Replace alloc_zeros::<half::bf16> → alloc_zeros::<f32>**

```bash
sed -i 's/alloc_zeros::<half::bf16>/alloc_zeros::<f32>/g' crates/ml/src/cuda_pipeline/gpu_weights.rs
```

- [ ] **Step 3: Delete CuriosityWeightSetBf16 and alloc_bf16_mirror**

These exist solely to maintain bf16 shadow copies for GemmEx. With cublasLtMatmul doing f32→TF32, the bf16 mirrors are dead:

Delete `struct CuriosityWeightSetBf16` (lines ~527-533).
Delete `fn alloc_bf16_mirror` (lines ~535-543).
Delete `pub fn convert_f32_to_bf16` (lines ~551-580+).

- [ ] **Step 4: Fix bf16_weight_ptrs_from_base → f32_weight_ptrs_from_base**

If the function `bf16_weight_ptrs_from_base` exists, rename it and update its pointer arithmetic. bf16 uses 2-byte stride, f32 uses 4-byte stride. The function computes offsets as `base + cumulative_size * sizeof(element)`:

```bash
grep -n 'bf16_weight_ptrs_from_base\|f32_weight_ptrs_from_base' crates/ml/src/cuda_pipeline/gpu_weights.rs
```

If it does byte-offset math with `* 2` (bf16 size), change to `* 4` (f32 size). If it uses element counts (not bytes), no change needed.

- [ ] **Step 5: Remove half crate imports**

```bash
sed -i '/use half/d' crates/ml/src/cuda_pipeline/gpu_weights.rs
sed -i '/half::bf16/d' crates/ml/src/cuda_pipeline/gpu_weights.rs  # only standalone lines
```

- [ ] **Step 6: Compile check**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -30`
Expected: Errors in modules that import weight structs. Fix any remaining type mismatches.

- [ ] **Step 7: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_weights.rs
git commit -m "refactor: gpu_weights all CudaSlice<half::bf16> → CudaSlice<f32>, delete bf16 mirror infrastructure"
```

---

### Task 5: Rewrite `gpu_dqn_trainer.rs` — Struct Fields + Constructor

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:500-3000`

The DQN trainer struct has ~50 `CudaSlice<half::bf16>` fields and a ~600-line constructor that allocates them all via `alloc_bf16()`. This task converts the struct definition and constructor. Methods are Task 6.

- [ ] **Step 1: Regex replace all CudaSlice<half::bf16> in struct fields**

```bash
sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```

This hits ~50 struct fields including:
- `spec_u_s1..spec_v_bg2` (spectral norm vectors)
- `params_bf16`, `target_params_bf16`
- `save_current_lp`, `save_projected` (C51 distributions)
- `per_sample_loss_buf`, `td_errors_buf`
- `q_out_buf`, `readback_buf`, `upload_staging_buf`, `scalars_readback_buf`
- `d_value_logits_bf16`, `d_adv_logits_bf16`
- `iqn_trunk_v`, `attn_bw_bf16_scratch`
- `causal_states_scratch`, `causal_q_scratch`
- `bn_hidden_buf`, `bn_concat_buf`

- [ ] **Step 2: Replace alloc_bf16() calls with alloc_f32()**

The private `alloc_bf16` function at the bottom of the file allocates `CudaSlice<half::bf16>`. Replace all calls:

```bash
sed -i 's/alloc_bf16(\([^)]*\))/alloc_f32(\1)/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```

Then rewrite the `alloc_bf16` function itself:

```rust
fn alloc_f32(
    stream: &Arc<CudaStream>,
    n: usize,
    label: &str,
) -> Result<CudaSlice<f32>, MLError> {
    stream.alloc_zeros::<f32>(n).map_err(|e| {
        MLError::ModelError(format!("alloc_f32 '{label}' ({n} elems): {e}"))
    })
}
```

Delete the old `alloc_bf16` function. Also replace any remaining `alloc_zeros::<half::bf16>` with `alloc_zeros::<f32>`.

- [ ] **Step 3: Fix size_of calculations**

Search for `size_of::<half::bf16>()` and replace with `size_of::<f32>()`:

```bash
sed -i 's/size_of::<half::bf16>()/size_of::<f32>()/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```

This affects VRAM budget calculations and memcpy byte counts. Critical: the VRAM budget will increase ~2x since f32 is 4 bytes vs bf16's 2 bytes. Verify total VRAM usage is within H100's 80GB budget. At batch=16384 with 26 weight tensors + activation saves, the increase is manageable (~2-4GB additional).

- [ ] **Step 4: Replace half::bf16::from_f32() and .to_f32()**

```bash
sed -i 's/half::bf16::from_f32(\([^)]*\))/\1/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
sed -i 's/half::bf16::ZERO/0.0_f32/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
sed -i 's/\.to_f32()//g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```

- [ ] **Step 5: Rename bf16-named fields**

Rename struct fields that have `bf16` in their name to reflect f32:

| Old Name | New Name |
|----------|----------|
| `params_bf16` | `params_flat` |
| `target_params_bf16` | `target_params_flat` |
| `d_value_logits_bf16` | `d_value_logits` |
| `d_adv_logits_bf16` | `d_adv_logits` |
| `attn_bw_bf16_scratch` | `attn_bw_scratch` |
| `aux_norm_bf16_scratch` | `aux_norm_scratch` (already done) |

Use replace_all for each:
```bash
sed -i 's/params_bf16/params_flat/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
sed -i 's/target_params_bf16/target_params_flat/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
sed -i 's/d_value_logits_bf16/d_value_logits/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
sed -i 's/d_adv_logits_bf16/d_adv_logits/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
sed -i 's/attn_bw_bf16_scratch/attn_bw_scratch/g' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```

- [ ] **Step 6: Delete f32_to_bf16_kernel and bf16_to_f32_kernel fields**

Remove these kernel function handles from the struct and constructor:
- `f32_to_bf16_kernel: CudaFunction` — DELETE field
- `bf16_to_f32_kernel: CudaFunction` — DELETE field
- `nan_check_bf16_kernel: CudaFunction` — KEEP but rename to reflect it now checks f32

Remove all launch sites that call these kernels (they were used to sync f32 master → bf16 shadow). With pure f32, no sync is needed.

- [ ] **Step 7: Delete the kernel loading for dead kernels**

In the `load_utility_kernels` function (~line 6340), remove the `f32_to_bf16` and `bf16_to_f32` kernel loads. Update the destructuring tuple.

- [ ] **Step 8: Compile check + fix remaining errors**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | wc -l`
Fix any remaining errors — likely in method signatures that reference old field names. Address one-by-one using compiler output.

- [ ] **Step 9: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "refactor: gpu_dqn_trainer struct + constructor — all CudaSlice<half::bf16> → CudaSlice<f32>, rename bf16 fields"
```

---

### Task 6: Rewrite `gpu_dqn_trainer.rs` — Methods

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs:3000-6600`

With the struct fields converted in Task 5, methods that use those fields need updating. This includes forward, backward, spectral norm, Adam, EMA sync, and diagnostic methods.

- [ ] **Step 1: Delete all f32_to_bf16_kernel launch sites**

Search and delete every block that launches `f32_to_bf16_kernel` (f32→bf16 shadow sync):

```bash
grep -n 'f32_to_bf16_kernel\|bf16_to_f32_kernel' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
```

These launch sites are approximately at lines: 992, 2209, 4063, 5566 (and their surrounding ~10 lines of launch_builder + arg chains). Delete each entire launch block. The purpose was to keep `params_bf16` in sync with `params_f32` — with pure f32, the master params ARE the training params.

- [ ] **Step 2: Fix spectral norm — now operates on f32 weights**

The spectral norm kernel (`spectral_norm_batched`) previously operated on bf16 weight pointers. Task 2 already converted the CUDA kernel to `float*`. Now fix the Rust side:

The descriptor array packs `[W_ptr, u_ptr, v_ptr, out_dim, in_dim, 0]` — the pointers now point to f32 params_flat instead of params_bf16. Update the descriptor-building code to use `self.params_flat.raw_ptr()` offset by f32 stride (4 bytes per element, not 2).

```bash
grep -n 'spec.*descriptor\|spectral.*desc\|spec_u_\|spec_v_' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -20
```

The critical fix: pointer offsets in the descriptor array. If the old code computed offsets as `base + cumulative_elements * 2` (bf16 bytes), change to `base + cumulative_elements * 4` (f32 bytes). If offsets were in elements (not bytes), no change needed — cudarc raw_ptr already gives byte addresses.

- [ ] **Step 3: Fix bw_d_h_s2_as_bf16 → bw_d_h_s2**

Rename the method and fix its return type:

```rust
pub fn bw_d_h_s2(&mut self) -> Result<&CudaSlice<f32>, MLError> {
    // ... copy d_h_s2 into attn_bw_scratch ...
    Ok(&self.attn_bw_scratch)
}
```

- [ ] **Step 4: Fix readback methods**

Methods that read bf16 back to host (via `dtoh_bf16_to_f32`) now read f32 directly:

```bash
grep -n 'dtoh_bf16_to_f32\|readback' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -20
```

Replace `dtoh_bf16_to_f32(&stream, &buf, &mut host)` with `dtoh_f32(&stream, &buf, &mut host)` (or the legacy alias works too since Task 3 made it forward to f32).

- [ ] **Step 5: Fix upload methods**

Methods that upload f32 host data → bf16 GPU (via `htod_f32_to_bf16`) now upload directly:

```bash
grep -n 'htod_f32_to_bf16\|upload' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | head -20
```

Replace `htod_f32_to_bf16(&stream, &data, &mut buf)` with `htod_f32(&stream, &data, &mut buf)`.

- [ ] **Step 6: Fix params accessor methods**

Methods like `params_bf16()`, `target_params_bf16()` need renaming:

```rust
pub fn params_flat(&self) -> &CudaSlice<f32> { &self.params_flat }
pub fn target_params_flat(&self) -> &CudaSlice<f32> { &self.target_params_flat }
```

- [ ] **Step 7: Fix memcpy byte calculations**

Any DtoD copy that computed bytes as `n * 2` (bf16 size) must change to `n * 4` (f32 size). Search:

```bash
grep -n '\* 2\b' crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | grep -i 'byte\|memcpy\|size\|len'
```

Or better: use `n * std::mem::size_of::<f32>()` everywhere.

- [ ] **Step 8: Compile check + fix remaining errors**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | wc -l`
Fix remaining errors iteratively. The compiler is the guide.

- [ ] **Step 9: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs
git commit -m "refactor: gpu_dqn_trainer methods — delete bf16 cast kernel launches, fix spectral norm + readback + upload for pure f32"
```

---

### Task 7: Rewrite `fused_training.rs`

**Files:**
- Modify: `crates/ml/src/trainers/dqn/fused_training.rs`

The fused training loop holds ensemble buffers, graph capture state, and method signatures that reference `CudaSlice<half::bf16>`.

- [ ] **Step 1: Regex replace all CudaSlice<half::bf16> → CudaSlice<f32>**

```bash
sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' crates/ml/src/trainers/dqn/fused_training.rs
```

This hits: `best_params_snapshot`, `ensemble_mean_q_buf`, `ensemble_var_q_buf`, `ensemble_diversity_loss_buf`, `ensemble_d_logits_buf`.

- [ ] **Step 2: Replace alloc_zeros::<half::bf16> → alloc_zeros::<f32>**

```bash
sed -i 's/alloc_zeros::<half::bf16>/alloc_zeros::<f32>/g' crates/ml/src/trainers/dqn/fused_training.rs
```

- [ ] **Step 3: Fix size_of and memset calls**

```bash
sed -i 's/size_of::<half::bf16>()/size_of::<f32>()/g' crates/ml/src/trainers/dqn/fused_training.rs
```

- [ ] **Step 4: Fix method signatures that return CudaSlice<half::bf16>**

Methods like `params_bf16_ptr()`, `target_params_bf16_ptr()`, and methods that return `&CudaSlice<half::bf16>` need updating to match the renamed fields from Task 5:

```bash
grep -n 'params_bf16\|bf16_ptr\|as_bf16' crates/ml/src/trainers/dqn/fused_training.rs
```

Rename to `params_flat_ptr()`, etc.

- [ ] **Step 5: Fix bw_d_h_s2_as_bf16 call**

The attention backward call at ~line 990:
```rust
let d_h_s2 = self.trainer.bw_d_h_s2()?;  // was bw_d_h_s2_as_bf16
attn.backward(d_h_s2, self.batch_size)?;
```

- [ ] **Step 6: Fix IQN→PER cast block (~line 1053)**

The comment says "cast f32 per_sample_loss → bf16 td_errors_buf". With pure f32, this is a direct f32 copy — no cast needed. Replace the cast kernel launch with a DtoD memcpy:

```rust
// IQN→PER loss: direct f32 copy (no bf16 cast needed)
unsafe {
    cudarc::driver::sys::cuMemcpyDtoDAsync_v2(
        self.trainer.td_errors_buf_ptr(),
        iqn_loss_ptr,
        batch_size * std::mem::size_of::<f32>(),
        self.stream.cu_stream(),
    );
}
```

- [ ] **Step 7: Update graph capture comments**

Remove references to "bf16→f32 sync" in graph capture documentation. The spectral norm graph no longer needs bf16 conversion.

- [ ] **Step 8: Remove half crate imports**

```bash
sed -i '/use half/d' crates/ml/src/trainers/dqn/fused_training.rs
```

- [ ] **Step 9: Compile check + fix**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -20`
Fix remaining errors.

- [ ] **Step 10: Commit**

```bash
git add crates/ml/src/trainers/dqn/fused_training.rs
git commit -m "refactor: fused_training — pure f32 ensemble buffers, delete bf16 cast in IQN→PER path"
```

---

### Task 8: Rewrite GPU Experience + Action + Walk-Forward + HER Modules

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_experience_collector.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_action_selector.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_walk_forward.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_her.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs`

These modules hold experience replay buffers, action selection buffers, and walk-forward dataset buffers — all `CudaSlice<half::bf16>`.

- [ ] **Step 1: Regex sweep all 5 files**

```bash
for f in gpu_experience_collector.rs gpu_action_selector.rs gpu_walk_forward.rs gpu_her.rs gpu_backtest_evaluator.rs; do
    sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/alloc_zeros::<half::bf16>/alloc_zeros::<f32>/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/size_of::<half::bf16>()/size_of::<f32>()/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/half::bf16::from_f32(\([^)]*\))/\1/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/half::bf16::ZERO/0.0_f32/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/\.to_f32()//g' "crates/ml/src/cuda_pipeline/$f"
done
```

- [ ] **Step 2: Fix gpu_experience_collector.rs — online_params_flat**

The field `online_params_flat` (was `online_params_flat` but typed `CudaSlice<half::bf16>`) holds a copy of trainer params for experience collection. The DtoD sync byte count must change from `total_params * 2` to `total_params * 4`:

```bash
grep -n 'total_weight_bytes\|total_params.*size_of\|memcpy.*params' crates/ml/src/cuda_pipeline/gpu_experience_collector.rs
```

Fix all byte-count calculations.

- [ ] **Step 3: Fix gpu_her.rs — HerBatch struct**

The `HerBatch` struct has `states`, `next_states`, `rewards`, `dones` as `CudaSlice<half::bf16>`. The regex already converted types; verify the `alloc_batch` helper also uses f32:

```bash
grep -n 'fn alloc.*batch\|fn new' crates/ml/src/cuda_pipeline/gpu_her.rs | head -5
```

- [ ] **Step 4: Fix gpu_walk_forward.rs — GpuDataset struct**

The `GpuDataset` struct holds the full walk-forward dataset on GPU. Same regex pattern, verify constructors.

- [ ] **Step 5: Remove half crate imports from all 5 files**

```bash
for f in gpu_experience_collector.rs gpu_action_selector.rs gpu_walk_forward.rs gpu_her.rs gpu_backtest_evaluator.rs; do
    sed -i '/use half/d' "crates/ml/src/cuda_pipeline/$f"
done
```

- [ ] **Step 6: Compile check + fix**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -30`
Fix remaining type mismatches.

- [ ] **Step 7: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_experience_collector.rs crates/ml/src/cuda_pipeline/gpu_action_selector.rs crates/ml/src/cuda_pipeline/gpu_walk_forward.rs crates/ml/src/cuda_pipeline/gpu_her.rs crates/ml/src/cuda_pipeline/gpu_backtest_evaluator.rs
git commit -m "refactor: experience + action + walk-forward + HER modules — pure f32 buffers"
```

---

### Task 9: Rewrite GPU Supporting Modules (Monitoring, Portfolio, Statistics, Training Guard, IQN, Attention)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_monitoring.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_portfolio.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_statistics.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_training_guard.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_iqn_head.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_attention.rs`
- Modify: `crates/ml/src/cuda_pipeline/signal_adapter.rs`
- Modify: `crates/ml/src/cuda_pipeline/batched_forward.rs`
- Modify: `crates/ml/src/cuda_pipeline/batched_backward.rs`

- [ ] **Step 1: Regex sweep all 9 files**

```bash
for f in gpu_monitoring.rs gpu_portfolio.rs gpu_statistics.rs gpu_training_guard.rs gpu_iqn_head.rs gpu_attention.rs signal_adapter.rs batched_forward.rs batched_backward.rs; do
    sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/alloc_zeros::<half::bf16>/alloc_zeros::<f32>/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/size_of::<half::bf16>()/size_of::<f32>()/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/half::bf16::from_f32(\([^)]*\))/\1/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/half::bf16::ZERO/0.0_f32/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/\.to_f32()//g' "crates/ml/src/cuda_pipeline/$f"
done
```

- [ ] **Step 2: Fix gpu_attention.rs — forward/backward signatures**

The attention module's `forward` takes `&CudaSlice<half::bf16>` and returns `&CudaSlice<half::bf16>`. After regex, verify the method signatures are correct and the internal memcpy byte counts use f32 size.

- [ ] **Step 3: Fix gpu_training_guard.rs — method signatures**

Training guard methods like `check_training_health` take `&CudaSlice<half::bf16>` Q-values. Verify types match after regex.

- [ ] **Step 4: Fix signal_adapter.rs — test functions**

Signal adapter has test functions that create bf16 data manually:
```rust
let bf16_uniform: Vec<half::bf16> = uniform.iter().map(|&v| half::bf16::from_f32(v)).collect();
```
After regex, these become identity operations. Simplify to direct f32 vec creation.

- [ ] **Step 5: Remove half crate imports from all files**

```bash
for f in gpu_monitoring.rs gpu_portfolio.rs gpu_statistics.rs gpu_training_guard.rs gpu_iqn_head.rs gpu_attention.rs signal_adapter.rs batched_forward.rs batched_backward.rs; do
    sed -i '/use half/d' "crates/ml/src/cuda_pipeline/$f"
done
```

- [ ] **Step 6: Compile check + fix**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -30`

- [ ] **Step 7: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_monitoring.rs crates/ml/src/cuda_pipeline/gpu_portfolio.rs crates/ml/src/cuda_pipeline/gpu_statistics.rs crates/ml/src/cuda_pipeline/gpu_training_guard.rs crates/ml/src/cuda_pipeline/gpu_iqn_head.rs crates/ml/src/cuda_pipeline/gpu_attention.rs crates/ml/src/cuda_pipeline/signal_adapter.rs crates/ml/src/cuda_pipeline/batched_forward.rs crates/ml/src/cuda_pipeline/batched_backward.rs
git commit -m "refactor: monitoring + portfolio + statistics + training guard + IQN + attention + signal adapter — pure f32"
```

---

### Task 10: Rewrite Non-DQN GPU Modules (PPO, IQL, Curiosity, Decision Transformer)

**Files:**
- Modify: `crates/ml/src/cuda_pipeline/gpu_ppo_collector.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs`
- Modify: `crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs`
- Modify: `crates/ml/src/cuda_pipeline/decision_transformer.rs`

- [ ] **Step 1: Regex sweep all 4 files**

```bash
for f in gpu_ppo_collector.rs gpu_iql_trainer.rs gpu_curiosity_trainer.rs decision_transformer.rs; do
    sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/alloc_zeros::<half::bf16>/alloc_zeros::<f32>/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/size_of::<half::bf16>()/size_of::<f32>()/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/half::bf16::from_f32(\([^)]*\))/\1/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/half::bf16::ZERO/0.0_f32/g' "crates/ml/src/cuda_pipeline/$f"
    sed -i 's/\.to_f32()//g' "crates/ml/src/cuda_pipeline/$f"
done
```

- [ ] **Step 2: Fix decision_transformer.rs — helper functions**

`raw_ptr_f32` and `raw_ptr_f32_mut` take `CudaSlice<half::bf16>` — after regex they take `CudaSlice<f32>`. Verify the `device_ptr` call still works with the new type.

- [ ] **Step 3: Fix gpu_curiosity_trainer.rs — adam_update helper**

The curiosity trainer has an inline Adam update that takes `&mut CudaSlice<half::bf16>` params and `&CudaSlice<half::bf16>` grads. After regex, verify types match and the kernel launch arg types are correct.

- [ ] **Step 4: Fix gpu_iql_trainer.rs — value head buffers**

IQL trainer has separate value head buffers (save_pre1, save_pre2, etc.) all typed bf16. After regex, verify constructor allocations match.

- [ ] **Step 5: Remove half crate imports**

```bash
for f in gpu_ppo_collector.rs gpu_iql_trainer.rs gpu_curiosity_trainer.rs decision_transformer.rs; do
    sed -i '/use half/d' "crates/ml/src/cuda_pipeline/$f"
done
```

- [ ] **Step 6: Compile check + fix**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -30`

- [ ] **Step 7: Commit**

```bash
git add crates/ml/src/cuda_pipeline/gpu_ppo_collector.rs crates/ml/src/cuda_pipeline/gpu_iql_trainer.rs crates/ml/src/cuda_pipeline/gpu_curiosity_trainer.rs crates/ml/src/cuda_pipeline/decision_transformer.rs
git commit -m "refactor: PPO + IQL + curiosity + decision transformer — pure f32 buffers"
```

---

### Task 11: Rewrite Non-Pipeline Rust Files (Trainers, Config, Smoke Tests)

**Files:**
- Modify: `crates/ml/src/trainers/dqn/trainer/mod.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/training_loop.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/metrics.rs`
- Modify: `crates/ml/src/trainers/dqn/trainer/action.rs`
- Modify: `crates/ml/src/trainers/dqn/config.rs`
- Modify: `crates/ml/src/trainers/dqn/smoke_tests/gpu_residency.rs`
- Modify: `crates/ml/src/trainers/dqn/smoke_tests/gradient_budget.rs`
- Modify: `crates/ml/src/trainers/dqn/smoke_tests/regression.rs`
- Modify: `crates/ml/src/trainers/dqn/smoke_tests/walk_forward.rs`
- Modify: `crates/ml/src/trainers/ppo.rs`
- Modify: `crates/ml/src/trainers/tlob.rs`
- Modify: `crates/ml/src/training_profile.rs`
- Modify: `crates/ml/src/hyperopt/adapters/mamba2.rs`
- Modify: `crates/ml/src/hyperopt/adapters/ppo.rs`

- [ ] **Step 1: Regex sweep all files**

```bash
find crates/ml/src/trainers crates/ml/src/hyperopt crates/ml/src/training_profile.rs -name '*.rs' | while read f; do
    sed -i 's/CudaSlice<half::bf16>/CudaSlice<f32>/g' "$f"
    sed -i 's/alloc_zeros::<half::bf16>/alloc_zeros::<f32>/g' "$f"
    sed -i 's/size_of::<half::bf16>()/size_of::<f32>()/g' "$f"
    sed -i 's/half::bf16::from_f32(\([^)]*\))/\1/g' "$f"
    sed -i 's/half::bf16::ZERO/0.0_f32/g' "$f"
    sed -i 's/\.to_f32()//g' "$f"
    sed -i 's/params_bf16/params_flat/g' "$f"
    sed -i 's/target_params_bf16/target_params_flat/g' "$f"
done
```

- [ ] **Step 2: Fix smoke test assertions**

Smoke tests may assert on bf16-precision values. With f32, tolerances can tighten:

```bash
grep -rn 'bf16\|bfloat' crates/ml/src/trainers/dqn/smoke_tests/ --include='*.rs'
```

Fix any remaining references to bf16 types or precision comments.

- [ ] **Step 3: Fix config.rs**

```bash
grep -rn 'bf16' crates/ml/src/trainers/dqn/config.rs
```

Remove any bf16-related config fields or comments.

- [ ] **Step 4: Remove half crate imports from all files**

```bash
find crates/ml/src/trainers crates/ml/src/hyperopt crates/ml/src/training_profile.rs -name '*.rs' -exec sed -i '/use half/d' {} \;
```

- [ ] **Step 5: Compile check + fix**

Run: `SQLX_OFFLINE=true cargo check -p ml --lib 2>&1 | grep 'error\[' | head -20`

- [ ] **Step 6: Commit**

```bash
git add crates/ml/src/trainers/ crates/ml/src/hyperopt/ crates/ml/src/training_profile.rs
git commit -m "refactor: trainers + config + smoke tests + hyperopt — pure f32, no half::bf16 references"
```

---

### Task 12: Dead Code Cleanup + Remove `half` Crate

**Files:**
- Modify: `crates/ml/Cargo.toml`
- Modify: `crates/ml/src/cuda_pipeline/mod.rs`
- Modify: `crates/ml/src/cuda_pipeline/common_device_functions.cuh`
- Modify: `crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs`
- Modify: Various files with dead imports

- [ ] **Step 1: Remove legacy bf16 function aliases from mod.rs**

Delete the deprecated aliases added in Task 3:
- `htod_f32_to_bf16` → DELETE (callers now use `htod_f32`)
- `clone_htod_f32_to_bf16` → DELETE
- `dtoh_bf16_to_f32` → DELETE
- `f32_slice_to_bf16` → DELETE
- `buf_bytes_bf16` → DELETE

Then update all remaining callers to use the new names:

```bash
grep -rn 'htod_f32_to_bf16\|clone_htod_f32_to_bf16\|dtoh_bf16_to_f32\|f32_slice_to_bf16\|buf_bytes_bf16' crates/ml/src/ --include='*.rs'
```

Replace each call site with the f32 equivalent.

- [ ] **Step 2: Rename bf16_ wrapper functions in common_device_functions.cuh**

The CUDA header still has `bf16_sqrt`, `bf16_log`, etc. as names (though they now operate on float). Decide: rename to `f32_*` OR delete entirely and inline the standard math calls. Recommended: delete the wrappers and use standard CUDA math directly. But this requires updating all call sites in 30 .cu files.

Alternative (lower risk): keep the wrapper names — they're just thin aliases to standard math. Add a comment: `/* Legacy names — these are pure f32 operations */`.

- [ ] **Step 3: Delete f32_to_bf16_kernel and bf16_to_f32_kernel from dqn_utility_kernels.cu**

These identity kernels (from Task 2) are now unused. Delete them.

- [ ] **Step 4: Delete dqn_nan_check_bf16 or rename to dqn_nan_check_f32**

If the kernel is still used, rename. If not, delete.

- [ ] **Step 5: Remove `half` crate from Cargo.toml**

```bash
grep -n 'half' crates/ml/Cargo.toml
```

If `half` is ONLY used for bf16 GPU buffers (not for f16 or other purposes), remove the dependency entirely. Check if any other crate in the workspace depends on half through ml:

```bash
grep -rn 'half' crates/*/Cargo.toml
```

If other crates use it for non-GPU purposes (e.g., f16 serialization), keep it in those crates but remove from ml.

- [ ] **Step 6: Verify zero bf16 references remain**

```bash
grep -rn 'bf16\|bfloat16\|half::bf16\|__nv_bfloat16' crates/ml/src/ --include='*.rs' --include='*.cu' | grep -v fxcache | grep -v '// ' | grep -v 'Legacy\|DEPRECATED'
```

Expected: 0 matches (excluding fxcache.rs disk format and comments).

- [ ] **Step 7: Full compile check**

Run: `SQLX_OFFLINE=true cargo check --workspace 2>&1 | grep 'error\[' | wc -l`
Expected: 0 errors.

- [ ] **Step 8: Commit**

```bash
git add -A
git commit -m "cleanup: remove half crate, delete all bf16 dead code — pure f32/TF32 pipeline complete"
```

---

### Task 13: Run Smoke Tests + Verify Correctness

**Files:**
- No modifications — verification only

- [ ] **Step 1: Run all DQN smoke tests**

```bash
SQLX_OFFLINE=true FOXHUNT_TEST_DATA=test_data/futures-baseline cargo test -p ml --lib -- smoke_tests --ignored --nocapture 2>&1 | tail -40
```

Expected: All 19 tests pass. Watch for:
- NaN in loss or Q-values (would indicate a buffer size/offset bug)
- CUDA launch errors (would indicate kernel arg mismatches)
- Assertion failures in gradient budget tests (thresholds may need loosening/tightening for f32)

- [ ] **Step 2: Run DQN unit tests**

```bash
SQLX_OFFLINE=true cargo test -p ml --lib -- dqn --nocapture 2>&1 | tail -20
```

- [ ] **Step 3: Verify VRAM usage is within budget**

Check that the f32 buffers fit in GPU memory. The VRAM increase from bf16→f32 is ~2x for weight/activation buffers. At batch=16384 on H100 (80GB), the budget should be fine. On RTX 3050 (4GB), smaller batch sizes might be needed for tests.

- [ ] **Step 4: Commit test results (if any test adjustments needed)**

If smoke test thresholds need adjustment for f32 precision, fix and commit:

```bash
git add crates/ml/src/trainers/dqn/smoke_tests/
git commit -m "test: adjust smoke test thresholds for pure f32 precision"
```
