[BUG] Operand initialized with a constant doesn't work with `inline_ptx` as an operand
Author: kainzhongCreated Sep 17, 2026Updated Sep 18, 2026
Labelsbug? - Needs TriageCuTe DSL
Which component has the problem?
CuTe DSL
Bug Report
Describe the bug
In CuTeDSL, if you define a variable as a constant like zero = Int32(0), inline_ptx lowers Int32(0) to the immediate operand 0. This causes ptxas to reject instructions whose operand grammar requires a register, such as max.xorsign.abs.bf16x2. There is currently no apparent way to request register materialization for a constant passed through read_only_args.
Steps/Code to reproduce bug
import cutlass
from cutlass import cute
from cutlass import Int32
@cute.kernel
def kernel():
value = cute.arch.thread_idx()[0]
zero = Int32(0)
cute.arch.inline_ptx(
"max.xorsign.abs.bf16x2 {$w0}, {$r0}, {$r1};",
write_only_types=[Int32],
read_only_args=[zero, value],
)
@cute.jit
def launch():
kernel().launch(grid=(1, 1, 1), block=(32, 1, 1))
if __name__ == "__main__":
cute.compile(launch)Outputs:
error: PTX assembly failed
--> /home/kainingz/GitHub/TransformerEngine/repro_cutedsl_inline_ptx_constant.py:21:0
|
19 |
20 |
> 21 | @cute.kernel
| ^
22 | def kernel():
23 | value = cute.arch.thread_idx()[0]
error: ptxas rejected the PTX generated for this kernel while compiling it to SASS.
note: target architecture: sm_100
note: ptxas log:
ptxas application ptx input, line 22; error : Arguments mismatch for instruction 'max'
ptxas fatal : Ptx assembly aborted due to errors
suggestion: read the ptxas log above; re-run with CUTE_DSL_KEEP=ptx and assemble the dumped PTX
with `ptxas -arch=sm_100 <dumped>.ptx` to iterate on the failureGenerated PTX:
//
// Generated by NVIDIA NVVM Compiler
//
// Compiler Build ID: CL-38501229
// Cuda compilation tools, release 13.4, V13.4.46
// Based on NVVM 23.0.0
//
.version 9.4
.target sm_100
.address_size 64
// .globl kernel_cutlass_kernel_0
.visible .entry kernel_cutlass_kernel_0()
.reqntid 32, 1, 1
{
.reg .b32 %r<3>;
mov.u32 %r2, %tid.x;
// begin inline asm
max.xorsign.abs.bf16x2 %r1, 0, %r2;
// end inline asm
ret;
}Expected behavior It should be
.reg .b32 %r0;
mov.b32 %r0, 0;
max.xorsign.abs.bf16x2 %r1, %r0, %r2;where zero is not a literal zero.
Environment details (please complete the following information): nvidia-cutlass-dsl 4.8.0.dev0
Additional context Add any other context about the problem here.
Source: NVIDIA/cutlass