biondizzle/nvfp4-megamoe-kernel
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

NVFP4-1.1: fix cute.arch.store signature - store(ptr, val) not store(ptr, val, dtype)

Browse Source
This commit is contained in:
biondizzle
2026-05-28 04:01:38 +00:00
parent a05a76bb6b
commit dc35d29811
2 changed files with 3 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
tests/unit/test_int32_cast.py
View File

@@ -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__":

4
tests/unit/test_nvfp4_1_1_quant.py
View File

@@ -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():