Zero initialize native NVFP4 block scales

This commit is contained in:
Daniel Maddern 2026-08-15 02:38:08 +07:00
parent 7f9b1560e4
commit fbfb3453de

View file

@ -294,7 +294,7 @@ std::vector<torch::Tensor> 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<<<scale_rows, threads, 0, stream>>>(