[QST]`print_tensor` cause "an illegal instruction was encountered"

Title: print_tensor triggers cudaErrorIllegalInstruction (715) on SM89 when multiple tensors are printed sequentially

PTENSOR(s2r_K_s);

Environment

• CUTLASS Version: 087c84df (Commit date: Mar 14, 2026)
• NVCC Version: 12.9
• GPU: NVIDIA GeForce RTX 4070 Ti (SM89)
• OS: Ubuntu 22.04 (Conda environment)

Description

I am implementing a FlashAttention-2 style kernel for learning purposes using CuTe. I encountered a very strange illegal memory access (which manifests as cudaErrorIllegalInstruction: 715 in compute-sanitizer) specifically when using print_tensor on multiple tensors within the same kernel.

The Symptom:

1. Defined a helper: #define PTENSOR(x) if(thread0()) {print(#x" "); print_tensor(x); print("\n");}
2. PTENSOR triggers illegal instruction (error 715), especially when placed inside a loop.

error code

  for(int ik = 0; ik < k_len; ++ik) {
    Tensor score = partition_fragment_C(tiled_mma, make_shape(Int<kBlockM>{}, Int<kBlockN>{}));
    Tensor score_RC = make_tensor(score.data(), convert_rowcol(score.layout()));
    clear(score);
    cp_async_wait<0>();
    __syncthreads();

    Tensor gV = local_tile(V, make_tile(Int<kBlockN>{}, Int<kHeadDim>{}), ik);
    auto g2s_gV_s = thr_g2s.partition_S(gV);
    copy(g2s, g2s_gV_s, g2s_sV_d);
    cp_async_fence();
    gemm_flash(score, rQ, rK, s2r_K_d, s2r_K_s, tiled_mma, s2r_K);
    if constexpr(IsCausal) {
        if(ik*kBlockN >= blockIdx.x*kBlockM) {
            causal_mask<kBlockN, kBlockM, kNWarps>(score_RC, ik);
        }
    }    
    PTENSOR(s2r_K_s);    // bug position
    // PTENSOR(s2r_Vt_s);

    if (ik < k_len-1) {
        gK = local_tile(K, make_tile(Int<kBlockN>{}, Int<kHeadDim>{}), ik+1);
        g2s_gK_s = thr_g2s.partition_S(gK);
        copy(g2s, g2s_gK_s, g2s_sK_d);
        cp_async_fence();
    }
    cp_async_wait<0>();
    __syncthreads();
    softmax_rescale(score_RC, score_max, score_sum, rOut, params.softmax_scale);
  }

Compute-Sanitizer Stack Trace

========= COMPUTE-SANITIZER
s2r_Q_s smem_ptr[16b](0x7fb449000000) o ((_8,_1),_1,(_2,_2)):((_1,_0),_0,(16,32)):
smem_ptr[16b](0x7fb449000000) o ((_8,_1),_1):((_1,_0),_0):
    0.00
    0.00
    0.00
    0.00
    0.00
    0.00
    0.00
    0.00
-========= Illegal instruction
=========     at cute::print(const char *)+0x6930 in print.hpp:152
=========     by thread (0,0,0) in block (0,0,0)
=========         Device Frame: void cute::print_tensor<cute::ViewEngine<cute::smem_ptr<cutlass::half_t *>>, cute::Layout<cute::tuple<cute::tuple<cute::C<(int)8>, cute::C<(int)1>>, cute::C<(int)1>, cute::tuple<cute::C<(int)2>, cute::C<(int)2>>>, cute::tuple<cute::tuple<cute::C<(int)1>, cute::C<(int)0>>, cute::C<(int)0>, cute::tuple<int, int>>>>(const cute::Tensor<T1, T2> &)+0x6720 in tensor.hpp:1013
=========         Device Frame: void flash::flash_atten2<Flash_fwd_kernel_traits<(int)64, (int)64, (int)64, (int)4, cutlass::half_t, Flash_kernel_traits<(int)64, (int)64, (int)64, (int)4, cutlass::half_t>>, Flash_fwd_params, (bool)1>(T2)+0x770 in flash_attention2.cu:203
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cuLaunchKernel [0x3dd1b4] in libcuda.so.1
=========         Host Frame:  [0x141f8] in libcudart.so.12
=========         Host Frame: cudaLaunchKernel [0x7d09d] in libcudart.so.12
=========         Host Frame: void flash::run_flash_fwd<Flash_fwd_kernel_traits<64, 64, 64, 4, cutlass::half_t, Flash_kernel_traits<64, 64, 64, 4, cutlass::half_t> >, true>(Flash_fwd_params&, CUstream_st*) [0x11d6c] in test
=========         Host Frame: flash::run_flash_attn2(Flash_fwd_params&) [0x5f73] in test
=========         Host Frame: flash_attention_v2(at::Tensor, at::Tensor, at::Tensor, bool, float) [0x16098] in test
=========         Host Frame: main [0x58ff] in test
========= 
========= Program hit cudaErrorIllegalInstruction (error 715) due to "an illegal instruction was encountered" on CUDA API call to cudaStreamSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: flash::run_flash_attn2(Flash_fwd_params&) [0x5cb6] in test
=========         Host Frame: flash_attention_v2(at::Tensor, at::Tensor, at::Tensor, bool, float) [0x16098] in test
=========         Host Frame: main [0x58ff] in test
========= 
========= Program hit cudaErrorIllegalInstruction (error 715) due to "an illegal instruction was encountered" on CUDA API call to cudaStreamSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: flash::run_flash_attn2(Flash_fwd_params&) [0x7ada] in test
=========         Host Frame: flash_attention_v2(at::Tensor, at::Tensor, at::Tensor, bool, float) [0x16098] in test
=========         Host Frame: main [0x58ff] in test
========= 
========= Program hit cudaErrorIllegalInstruction (error 715) due to "an illegal instruction was encountered" on CUDA API call to cudaStreamSynchronize.
=========     Saved host backtrace up to driver entry point at error
=========         Host Frame: flash::run_flash_attn2(Flash_fwd_params&) [0x7aeb] in test
=========         Host Frame: flash_attention_v2(at::Tensor, at::Tensor, at::Tensor, bool, float) [0x16098] in test
=========         Host Frame: main [0x58ff] in test
========= 
Cuda failure 715 flash_attention2.cu:422: 'an illegal instruction was encountered'
========= Target application returned an error
========= ERROR SUMMARY: 4 errors

Reproducible Code

git@github.com:happyflathead/print-tensor-bug.git

Start with the reproducible code in the linked print-tensor-bug repository and reproduce the SM89 failure using the shown PTENSOR(s2r_K_s) placement. Inspect cute::print in print.hpp:152, cute::print_tensor in tensor.hpp:1013, and the kernel call in flash_attention2.cu:203, comparing sequential and looped tensor printing. Done means the illegal-instruction behavior is isolated and its cause or a confirmed minimal reproduction is documented.