From 2ae5ab31940fc5af66bbdcbe68f91e04cc8c6be5 Mon Sep 17 00:00:00 2001 From: jgrusewski Date: Sun, 19 Apr 2026 01:38:51 +0200 Subject: [PATCH] cleanup: remove all stale bf16/BF16 references from comments MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Pure f32/TF32 pipeline — no bfloat16 anywhere. Co-Authored-By: Claude Opus 4.6 (1M context) --- crates/ml/src/cuda_pipeline/batched_forward.rs | 6 +++--- crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs | 14 +++++++------- .../src/cuda_pipeline/gpu_experience_collector.rs | 2 +- crates/ml/src/cuda_pipeline/mod.rs | 2 +- 4 files changed, 12 insertions(+), 12 deletions(-) diff --git a/crates/ml/src/cuda_pipeline/batched_forward.rs b/crates/ml/src/cuda_pipeline/batched_forward.rs index d99857ca7..1d38c781f 100644 --- a/crates/ml/src/cuda_pipeline/batched_forward.rs +++ b/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, diff --git a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs b/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs index 271d18693..e265904e7 100644 --- a/crates/ml/src/cuda_pipeline/gpu_dqn_trainer.rs +++ b/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, /// Gradient w.r.t. branch logits: [B, (B0+B1+B2+B3)*NA] — f32 for native atomicAdd d_adv_logits_buf: CudaSlice, - /// 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, - /// 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, @@ -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::(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::()) 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 diff --git a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs b/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs index fc931aec3..371ba1946 100644 --- a/crates/ml/src/cuda_pipeline/gpu_experience_collector.rs +++ b/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 diff --git a/crates/ml/src/cuda_pipeline/mod.rs b/crates/ml/src/cuda_pipeline/mod.rs index 2c80495cb..dc9f29daa 100644 --- a/crates/ml/src/cuda_pipeline/mod.rs +++ b/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