cleanup: remove all stale bf16/BF16 references from comments

Pure f32/TF32 pipeline — no bfloat16 anywhere.

Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

crates/ml/src/cuda_pipeline/batched_forward.rs

@@ -1,6 +1,6 @@
 #![allow(unsafe_code)]
 
-//! cuBLAS BF16 batched forward pass for the DQN trainer.
+//! cuBLAS batched forward pass for the DQN trainer (f32/TF32).
 //!
 //! Replaces the 1-warp-per-sample fused kernel with cuBLAS matrix multiplications
 //! that process the entire batch in a single GEMM call per layer. This raises GPU
@@ -1846,9 +1846,9 @@ fn create_cached_fwd_gemm_desc_relu_bias(
     }
 }
 
-// ── Compute BF16 weight pointers from flat params_buf ───────────────────────
+// ── Compute weight pointers from flat params_buf ────────────────────────────
 
-/// Compute the 20 raw BF16 device pointers into a flat params_buf at GOFF_* offsets.
+/// Compute the 20 raw device pointers into a flat params_buf at GOFF_* offsets.
 ///
 /// The flat buffer layout matches `compute_param_sizes()`:
 /// [w_s1, b_s1, w_s2, b_s2, w_v1, b_v1, w_v2, b_v2,

crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs

@@ -255,7 +255,7 @@ impl Default for GpuDqnTrainConfig {
             lr: 3e-5,     // Conservative default for stable training
             beta1: 0.9,
             beta2: 0.999,
-            epsilon: 1e-3, // BF16-safe
+            epsilon: 1e-3,
             weight_decay: 1e-4,  // Standard L2 regularization strength
             max_grad_norm: 10000.0, // Safety-only — per-component clip removed, raw norms ~4000
             spectral_norm_sigma_max: 3.0,
@@ -1177,9 +1177,9 @@ pub struct GpuDqnTrainer {
     d_value_logits_buf: CudaSlice<f32>,
     /// Gradient w.r.t. branch logits: [B, (B0+B1+B2+B3)*NA] — f32 for native atomicAdd
     d_adv_logits_buf: CudaSlice<f32>,
-    /// BF16 staging for backward pass: cast from f32 d_value_logits before cuBLAS GEMM
+    /// Staging for backward pass value logits gradient
     d_value_logits: CudaSlice<f32>,
-    /// BF16 staging for backward pass: cast from f32 d_adv_logits before cuBLAS GEMM
+    /// Staging for backward pass advantage logits gradient
     d_adv_logits: CudaSlice<f32>,
 
 
@@ -4781,7 +4781,7 @@ impl GpuDqnTrainer {
             .map_err(|e| MLError::ModelError(format!("alloc d_value_logits_mse f32: {e}")))?;
         let d_adv_logits_mse = stream.alloc_zeros::<f32>(b * total_branch_atoms + 32 * 4)
             .map_err(|e| MLError::ModelError(format!("alloc d_adv_logits_mse f32: {e}")))?;
-        // BF16 staging buffers — cast from f32 before cuBLAS backward GEMM
+        // Staging buffers for backward pass cuBLAS GEMM
         let d_value_logits = alloc_f32(&stream, b * pad32(config.num_atoms), "d_value_logits")?;
         let d_adv_logits = alloc_f32(&stream, b * total_branch_atoms + 32 * 4, "d_adv_logits")?;
 
@@ -8513,7 +8513,7 @@ impl GpuDqnTrainer {
     /// 1-warp-per-sample shared-memory-bound kernel with cuBLAS DGEMM which
     /// uses tensor cores and efficiently tiles across the full SM.
     ///
-    /// Returns `Ok(false)` if cuBLAS is not initialized (caller falls back to BF16 kernel).
+    /// Returns `Ok(false)` if cuBLAS is not initialized.
     pub(crate) fn launch_cublas_forward(&self) -> Result<(), MLError> {
         let cublas = &self.cublas_forward;
 
@@ -9087,7 +9087,7 @@ impl GpuDqnTrainer {
         let b3_i32 = b3 as i32;
 
         // Shared memory: float arrays (support + val + adv + lp + proj + current_lp + reduce)
-        // Float (4 bytes/elem) for numerically stable softmax — BF16 exp() overflows at logit > 11.1
+        // Float (4 bytes/elem) for numerically stable softmax
         let max_branch = b0.max(b1).max(b2).max(b3);
         let shmem_floats = na + na + max_branch * na + na + na + na + 8;
         let shmem_bytes = (shmem_floats * std::mem::size_of::<f32>()) as u32;
@@ -10768,7 +10768,7 @@ fn pad32(n: usize) -> usize {
 
 /// Round up to the next multiple of 128 (CUTLASS K-tile alignment).
 ///
-/// cuBLAS CUTLASS kernels use 128-element K-tiles for BF16 GemmEx.
+/// cuBLAS CUTLASS kernels use 128-element K-tiles for TF32 GemmEx.
 /// When the K dimension (state_dim) is not a multiple of 128, CUTLASS
 /// reads past the row boundary of the B-matrix.  Padding each row of
 /// the states buffer to `pad128(state_dim)` with zeros eliminates OOB

crates/ml/src/cuda_pipeline/gpu_experience_collector.rs

@@ -1760,7 +1760,7 @@ impl GpuExperienceCollector {
             unsafe {
                 self.stream
                     .launch_builder(&self.hindsight_relabel_kernel)
-                    .arg(targets_buf)         // const bfloat16* targets
+                    .arg(targets_buf)         // const float* targets
                     .arg(&actions)            // const int* actions
                     .arg(&bar_indices_gpu)    // const int* bar_indices
                     .arg(&rewards)            // float* rewards

crates/ml/src/cuda_pipeline/mod.rs

@@ -46,7 +46,7 @@ const MAX_UPLOAD_BYTES: usize = 2 * 1024 * 1024 * 1024;
 /// Round dimension up to nearest multiple of 8 for tensor core alignment.
 ///
 /// Tensor cores on Ampere (A100, L4, L40S) and Hopper (H100) require
-/// matrix dimensions divisible by 8 (FP16/BF16) for optimal throughput.
+/// matrix dimensions divisible by 8 for optimal TF32 throughput.
 /// This pads dimensions like state_dim=45 → 48 to enable tensor core paths.
 pub fn align_to_tensor_cores(dim: usize) -> usize {
     (dim + 7) & !7