From fbfb3453de1c0766641f90cea49a0830b06b4ba7 Mon Sep 17 00:00:00 2001 From: Daniel Maddern Date: Sat, 15 Aug 2026 02:38:08 +0700 Subject: [PATCH] Zero initialize native NVFP4 block scales --- src/h3_blackwell_runtime/csrc/nvfp4_scale.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu b/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu index 5847167..fd4e29e 100644 --- a/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu +++ b/src/h3_blackwell_runtime/csrc/nvfp4_scale.cu @@ -294,7 +294,7 @@ std::vector quantize_nvfp4_bf16_cuda(torch::Tensor input, torch:: const int64_t scale_rows = roundup(q_rows, 128); const int64_t scale_cols = roundup(q_cols / 8, 4); auto qdata = torch::empty({q_rows, q_cols}, input.options().dtype(torch::kUInt8)); - auto block_scale = torch::empty({scale_rows, scale_cols}, input.options().dtype(torch::kFloat8_e4m3fn)); + auto block_scale = torch::zeros({scale_rows, scale_cols}, input.options().dtype(torch::kUInt8)).view(torch::kFloat8_e4m3fn); auto stream = at::cuda::getCurrentCUDAStream(); constexpr int threads = 256; quantize_nvfp4_bf16_kernel<<>>(