[BUG] CuTe DSL: SIGSEGV compiling integer ** (int32/int64), including 2 ** 3 #3735

Which component has the problem?

CuTe DSL

Bug Report

Describe the bug

Compiling a CUDA kernel containing x ** y, where both operands are loaded
from integer CUDA Tensors, terminates the Python process with SIGSEGV (exit
status 139) inside cute.compile().

The reproducer uses only small positive inputs: x = 2, y = 3, with expected
output 8. It does not involve zero bases, negative exponents, overflow, or
Python-valued constant evaluation. Compilation does not return, and the kernel
is not launched.

This reproduces for both int32 and int64 with native NVIDIA CuTeDSL wheels:

• 4.5.2;
• 4.8.0 (final release, cu12);
• 4.8.0.dev0 (cu12), tested separately before checking the final release.

For 4.8.0 final, the frontend and compiler were freshly installed into an
isolated target from matching 4.8.0 wheels. No source changes or custom compiler
forks were used.

Steps/Code to reproduce bug

Save this as minimal_integer_pow.py in an environment with native NVIDIA
CuTeDSL and CUDA-enabled PyTorch:

import argparse

import cutlass
import cutlass.cute as cute
import torch
from cutlass.cute.runtime import from_dlpack


@cute.kernel
def pow_kernel(a: cute.Tensor, b: cute.Tensor, out: cute.Tensor):
    i, _, _ = cute.arch.thread_idx()
    x = a[i]
    y = b[i]
    out[i] = x ** y


@cute.jit
def launch(a: cute.Tensor, b: cute.Tensor, out: cute.Tensor):
    pow_kernel(a, b, out).launch(grid=(1, 1, 1), block=(128, 1, 1))


def main():
    parser = argparse.ArgumentParser()
    parser.add_argument("--dtype", choices=("int32", "int64"), default="int32")
    args = parser.parse_args()
    print("CuTeDSL", cutlass.__version__, flush=True)
    print("PyTorch", torch.__version__, "CUDA runtime", torch.version.cuda,
          flush=True)
    print("GPU", torch.cuda.get_device_name(0), flush=True)
    print("dtype", args.dtype, "runtime inputs: a=2, b=3", flush=True)
    dtype = getattr(torch, args.dtype)
    a = torch.full((128,), 2, dtype=dtype, device="cuda")
    b = torch.full((128,), 3, dtype=dtype, device="cuda")
    out = torch.full_like(a, -123)
    handles = [from_dlpack(t, assumed_align=16) for t in (a, b, out)]
    print("COMPILE_BEGIN", flush=True)
    compiled = cute.compile(launch, *handles)
    print("COMPILE_OK", flush=True)
    compiled(*handles)
    torch.cuda.synchronize()
    result = out.cpu()
    assert torch.equal(result, torch.full_like(result, 8)), result
    print("PASS: synchronized output = 8", flush=True)


if __name__ == "__main__":
    main()

Run each dtype in a separate process:

python minimal_integer_pow.py --dtype int32
echo "exit=$?"
python minimal_integer_pow.py --dtype int64
echo "exit=$?"

Observed with 4.8.0 final:

CuTeDSL 4.8.0
PyTorch 2.7.0+cu128 CUDA runtime 12.8
GPU NVIDIA A100-PCIE-40GB
dtype int32 runtime inputs: a=2, b=3
COMPILE_BEGIN
PROCESS_EXIT_STATUS 139

COMPILE_OK is never reached. The shell reports Segmentation fault.
The int64 case also exits with status 139 during compilation.
4.5.2 additionally prints !!!!!!! Segfault encountered !!!!!!!.

Expected behavior

The kernel should compile and produce 8 for all elements. If runtime integer
exponentiation is not supported by this compilation path, an explicit compiler
diagnostic would be preferable to a process-level crash.

Environment details (please complete the following information):

• Environment location: Docker, Linux x86_64.
• GPU: NVIDIA A100-PCIE-40GB (SM80).
• Python: 3.10.
• PyTorch: 2.7.0+cu128; PyTorch CUDA runtime: 12.8.
• CUDA Toolkit: 12.9.
• NVIDIA driver: 610.43.02.
• CuTeDSL: 4.5.2; 4.8.0 final with nvidia-cutlass-dsl, libs-base,
  libs-core, and libs-cu12 all at 4.8.0; separately 4.8.0.dev0 cu12.
• Reproduced with file caching disabled (CUTE_DSL_ENABLE_FILE_CACHING=0).
• The reproducer compiles an @cute.jit launcher, not a bare @cute.kernel.

The container originally had older user-space driver libraries. Matching host
driver libraries were selected for these test processes without replacing the
system installation. The CUDA execution controls below successfully compiled,
launched, synchronized, and checked device output with that same setup.

Additional context

Control cases use the same Tensor-loaded inputs 2 and 3, the same
128-thread block, and the same DLPack/compilation bridge. Successful cases
synchronize the device and check every output, rather than only compiling.

Integer true division controls use float32 output. These are bounded controls,
not exhaustive coverage of all scalar operations or input domains.

The 4.5.2 pre-lowering IR contains math.ipowi on the loaded integer values.
An earlier GDB run stops inside the native NVIDIA _cutlass_ir.so, but does
not provide useful symbols. We have not identified a specific internal pass
root cause or established that every affected version has the same root cause.

We found the older floating-point power discussion #2296, but it concerns
v ** 1.3, increased compilation time, and a suspected runtime hang, rather
than this positive-integer compilation SIGSEGV. We could not find a matching
public issue or a documented fix for this integer-power reproducer.

Could you help clarify whether Tensor-loaded integer ** is intended to be
supported, and whether this crash matches a known issue or existing fix?