From dc35d2981140e28747596b9ff27c76125626240b Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 04:01:38 +0000 Subject: [PATCH] NVFP4-1.1: fix cute.arch.store signature - store(ptr, val) not store(ptr, val, dtype) --- tests/unit/test_int32_cast.py | 2 +- tests/unit/test_nvfp4_1_1_quant.py | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/unit/test_int32_cast.py b/tests/unit/test_int32_cast.py index a2d5a73b..5521b054 100644 --- a/tests/unit/test_int32_cast.py +++ b/tests/unit/test_int32_cast.py @@ -14,7 +14,7 @@ def test_int32_cast_kernel( if tidx == cutlass.Int32(0): f = cute.arch.load(input_f32.iterator, cutlass.Float32) i = cutlass.Int32(f) # float-to-int conversion - cute.arch.store(output_i32.iterator, i, cutlass.Int32) + cute.arch.store(output_i32.iterator, i) if __name__ == "__main__": diff --git a/tests/unit/test_nvfp4_1_1_quant.py b/tests/unit/test_nvfp4_1_1_quant.py index 6ad29469..7591272f 100644 --- a/tests/unit/test_nvfp4_1_1_quant.py +++ b/tests/unit/test_nvfp4_1_1_quant.py @@ -61,7 +61,7 @@ def fp4_quant_test_kernel( bs_dequant = fp8_e4m3_to_float32(sf_bits) # Write FP8 scale byte (Int32 holding uint8) - cute.arch.store(out_sf.iterator, sf_bits, cutlass.Int32) + cute.arch.store(out_sf.iterator, sf_bits) # ── Pass 2: Quantize and pack ── for i in cutlass.range(8, unroll=1): @@ -80,7 +80,7 @@ def fp4_quant_test_kernel( packed = (nibble1 << cutlass.Int32(4)) | nibble0 # Write packed byte as Int32 - cute.arch.store(out_fp4.iterator + i * cutlass.Int32(4), packed, cutlass.Int32) + cute.arch.store(out_fp4.iterator + i * cutlass.Int32(4), packed) def run_test():