From 2087eaef49cf16310d36139cf83f771898a88402 Mon Sep 17 00:00:00 2001
From: biondizzle <biondizzle@gmail.com>
Date: Thu, 28 May 2026 04:33:38 +0000
Subject: [PATCH] NVFP4-1.1: minimal threshold rounding test

---
 tests/unit/test_threshold_round.py | 31 ++++++++++++++++++++++++++++++
 1 file changed, 31 insertions(+)
 create mode 100644 tests/unit/test_threshold_round.py

diff --git a/tests/unit/test_threshold_round.py b/tests/unit/test_threshold_round.py
new file mode 100644
index 00000000..64cd6f69
--- /dev/null
+++ b/tests/unit/test_threshold_round.py
@@ -0,0 +1,31 @@
+"""Minimal test: just the threshold rounding function in CuTeDSL."""
+import torch
+import cutlass
+import cutlass.cute as cute
+import cutlass.torch as cutlass_torch
+from dsv4.kernels.gemm.fp4_quant import round_rne_u0_8
+
+
+@cute.kernel
+def threshold_test_kernel(
+    input_f32: cute.Tensor,    # (1,) Float32 input
+    output_i32: cute.Tensor,    # (1,) Int32 output
+):
+    tidx, _, _ = cute.arch.thread_idx()
+    if tidx == cutlass.Int32(0):
+        x = cute.arch.load(input_f32.iterator, cutlass.Float32)
+        result = round_rne_u0_8(x)
+        cute.arch.store(output_i32.iterator, result)
+
+
+if __name__ == "__main__":
+    x = torch.tensor([3.7], dtype=torch.float32, device='cuda')
+    out = torch.zeros(1, dtype=torch.int32, device='cuda')
+    xc = cutlass_torch.from_dlpack(x).mark_layout_dynamic(leading_dim=0)
+    oc = cutlass_torch.from_dlpack(out).mark_layout_dynamic(leading_dim=0)
+    
+    print("Compiling...")
+    compiled = cute.compile(threshold_test_kernel, xc, oc)
+    print("Running...")
+    compiled(xc, oc)
+    print(f'round(3.7) = {out.item()} (expected 4)')