From 08a65e287e55099255ce178c5b5f58ce42dfd461 Mon Sep 17 00:00:00 2001
From: "Kevin H. Luu" <khluu000@gmail.com>
Date: Sat, 17 Oct 2026 00:52:17 -0700
Subject: [PATCH] [Bugfix][Kernel] hadacore: don't load past the input on a
 partial last warp (#61944)

Signed-off-by: khluu <khluu000@gmail.com>
Co-authored-by: Claude Opus 5.5 <noreply@anthropic.com>
Co-authored-by: Jee Jee Li <pandaleefree@gmail.com>
Co-authored-by: vllm-agent <claw@inferact.ai>
---
 .../hadamard/hadacore/hadamard_transform_cuda.cu           | 7 +++++++
 1 file changed, 7 insertions(+)

diff --git a/csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu b/csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu
index 2575b20b8e..e86c518204 100644
--- a/csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu
+++ b/csrc/quantization/hadamard/hadacore/hadamard_transform_cuda.cu
@@ -135,6 +135,13 @@ hadamard_transform_kernel(b16* a, b16* out, int total_num_chunks) {
 
     #pragma unroll
     for (int64_t k = 0; k < num_chunks; k++) {
+        // The last warp may own fewer chunks than num_chunks: loading the
+        // rest would read past the end of `a`. The cp.async.wait_group counts
+        // below already assume only real_num_chunks groups were committed.
+        if constexpr(enable_mask) {
+            if (k >= real_num_chunks)
+                break;
+        }
         size_t shared_ptr = __cvta_generic_to_shared(b_frag_ptr);
         #if (__CUDA_ARCH__ >= 900) // SM90
             asm volatile(