[BUG] Operand initialized with a constant doesn't work with `inline_ptx` as an operand

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 failure

Generated 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.

Start by running the provided CuTe DSL reproduction and preserve the generated PTX with CUTE_DSL_KEEP=ptx. Inspect how inline_ptx handles read_only_args when compiling Int32(0), then compare the output with the expected register-materialized form and re-run ptxas to confirm the instruction is accepted.