From 4806e9ba117552c5a2a370ac39ef5474288305ee Mon Sep 17 00:00:00 2001 From: biondizzle Date: Thu, 28 May 2026 04:49:02 +0000 Subject: [PATCH] test: llvm.inline_asm with Int32._mlir_type matching cvt_i8_bf16 pattern --- tests/unit/test_ptx_llvm.py | 90 ++++++++++++++++++++++++++++++ tests/unit/test_ptx_llvm_runner.py | 19 +++++++ 2 files changed, 109 insertions(+) create mode 100644 tests/unit/test_ptx_llvm.py create mode 100644 tests/unit/test_ptx_llvm_runner.py diff --git a/tests/unit/test_ptx_llvm.py b/tests/unit/test_ptx_llvm.py new file mode 100644 index 00000000..35055bda --- /dev/null +++ b/tests/unit/test_ptx_llvm.py @@ -0,0 +1,90 @@ +"""Test: try llvm.inline_asm matching cvt_i8_bf16 pattern exactly.""" +import torch +import cutlass.cute as cute +import cutlass.torch as cutlass_torch +from cutlass.cutlass_dsl import dsl_user_op +from cutlass._mlir.dialects import llvm +from cutlass.cute.typing import Float32, Int32 +import cutlass + + +# Approach 1: llvm.inline_asm with Int32._mlir_type (matching cvt_i8_bf16 pattern) +@dsl_user_op +def f32_to_i32_rni_v1(x: Float32, *, loc=None, ip=None) -> Int32: + val_i32 = llvm.inline_asm( + Int32._mlir_type(), + [Float32(x).ir_value(loc=loc, ip=ip)], + "cvt.rni.s32.f32 $0, $1;", + "=r,f", + has_side_effects=False, + is_align_stack=False, + asm_dialect=llvm.AsmDialect.AD_ATT, + loc=loc, + ip=ip, + ) + return Int32(val_i32) + + +# Approach 2: llvm.inline_asm without asm_dialect (default) +@dsl_user_op +def f32_to_i32_rni_v2(x: Float32, *, loc=None, ip=None) -> Int32: + val_i32 = llvm.inline_asm( + Int32._mlir_type(), + [Float32(x).ir_value(loc=loc, ip=ip)], + "cvt.rni.s32.f32 $0, $1;", + "=r,f", + has_side_effects=False, + is_align_stack=False, + loc=loc, + ip=ip, + ) + return Int32(val_i32) + + +# Approach 3: Using multi-line block like cvt_i8_bf16 +@dsl_user_op +def f32_to_i32_rni_v3(x: Float32, *, loc=None, ip=None) -> Int32: + val_i32 = llvm.inline_asm( + Int32._mlir_type(), + [Float32(x).ir_value(loc=loc, ip=ip)], + "{\n\tcvt.rni.s32.f32 $0, $1;\n\t}", + "=r,f", + has_side_effects=False, + is_align_stack=False, + loc=loc, + ip=ip, + ) + return Int32(val_i32) + + +KERNELS = { + "v1_mlir_type": (f32_to_i32_rni_v1, None), + "v2_no_asm_dialect": (f32_to_i32_rni_v2, None), + "v3_multiline": (f32_to_i32_rni_v3, None), +} + +import sys +approach = sys.argv[1] if len(sys.argv) > 1 else "v1_mlir_type" +func = KERNELS[approach][0] + + +@cute.kernel +def test_k(inp: cute.Tensor, out: cute.Tensor): + tidx, _, _ = cute.arch.thread_idx() + if tidx == Int32(0): + x = cute.arch.load(inp.iterator, Float32) + r = func(x) + cute.arch.store(out.iterator, r) + + +if __name__ == "__main__": + x = torch.tensor([3.7], dtype=torch.float32, device='cuda') + o = 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(o).mark_layout_dynamic(leading_dim=0) + print(f"Approach: {approach}") + print("Compiling...") + compiled = cute.compile(test_k, xc, oc) + print("Running...") + compiled(xc, oc) + print(f"Result: {o.item()} (expected 4)") diff --git a/tests/unit/test_ptx_llvm_runner.py b/tests/unit/test_ptx_llvm_runner.py new file mode 100644 index 00000000..1766bf49 --- /dev/null +++ b/tests/unit/test_ptx_llvm_runner.py @@ -0,0 +1,19 @@ +"""Runner: test each llvm.inline_asm approach in separate process.""" +import subprocess +import sys + +for approach in ["v1_mlir_type", "v2_no_asm_dialect", "v3_multiline"]: + print(f"\n{'='*50}") + print(f"Testing: {approach}") + print(f"{'='*50}") + result = subprocess.run( + [sys.executable, "tests/unit/test_ptx_llvm.py", approach], + capture_output=True, + text=True, + timeout=120, + ) + print(result.stdout) + if result.stderr: + # Only show last 500 chars of stderr (LLVM ERROR is usually at the end) + print(f"STDERR (last 500): ...{result.stderr[-500:]}") + print(f"Exit code: {result.returncode}")