Looking for hints/suggestions for debugging TMA-related crash

[Artem-B]

I could use some advice on how to debug a TMA-related issue.

I'm debugging a handful of tests that appear to be crashing on SM90_TMA_LOAD_MULTICAST_3D::copy() in cutlass/include/cute/arch/copy_sm90_tma.hpp:

cutlass/include/cute/arch/copy_sm90_tma.hpp

Lines 631 to 643 in b78588d

	copy(void const* desc_ptr, uint64_t* mbar_ptr, uint16_t multicast_mask, uint64_t cache_hint,
	void * smem_ptr,
	int32_t const& crd0, int32_t const& crd1, int32_t const& crd2)
	{
	#if defined(CUTE_ARCH_TMA_SM90_ENABLED)
	uint64_t gmem_int_desc = reinterpret_cast<uint64_t>(desc_ptr);
	uint32_t smem_int_mbar = cast_smem_ptr_to_uint(mbar_ptr);
	uint32_t smem_int_ptr = cast_smem_ptr_to_uint(smem_ptr);
	cutlass::arch::synclog_emit_tma_load(__LINE__, gmem_int_desc, smem_int_mbar, smem_int_ptr);
	asm volatile (
	"cp.async.bulk.tensor.3d.shared::cluster.global.mbarrier::complete_tx::bytes.multicast::cluster.L2::cache_hint"
	" [%0], [%1, {%4, %5, %6}], [%2], %3, %7;"
	:

The problem is that I'm utterly failing at figuring out what makes the kernel crash with an "Illegal instruction".

CUDA Exception: Warp Illegal Instruction
(cuda-gdb) x/10i $pc-0x40
   0x7ffdc5b900e0:      @P3 R2UR UR9, R49
   0x7ffdc5b900f0:      @P3 R2UR UR8, R52
   0x7ffdc5b90100:      @P3 PLOP3.LUT P4, PT, P3, PT, PT, 0x8, 0x0
*> 0x7ffdc5b90110:      UTMALDG.3D.MULTICAST [UR8], [UR14], UR4, desc[UR6]
=> 0x7ffdc5b90120:      PLOP3.LUT P3, PT, PT, PT, PT, 0x8, 0x0

(cuda-gdb) p/x $UR8
$1 = 0x500
(cuda-gdb) p/x $UR14
$2 = 0xedfff5c0
(cuda-gdb) p/x $UR15
$3 = 0x7ffd
(cuda-gdb) x/32x 0x7ffdedfff5c0
0x7ffdedfff5c0: 0xb5e00000      0x00007ffd      0x000563a0      0x00000001
0x7ffdedfff5d0: 0x00000000      0x00000000      0x00000000      0x00000000
0x7ffdedfff5e0: 0x00000003      0x00000003      0x00000000      0x00000000
0x7ffdedfff5f0: 0x00000000      0x1f000000      0x0000001f      0x00000000
0x7ffdedfff600: 0x00001000      0x00000000      0x00000400      0x00000000
0x7ffdedfff610: 0x00000000      0x00000000      0x00000000      0x00000000
0x7ffdedfff620: 0x00000000      0x00000000      0x00000000      0x00000000
0x7ffdedfff630: 0x00000000      0x00000000      0x00000000      0x00000000
(cuda-gdb) p/x $UR4
$4 = 0x30000
(cuda-gdb) p/x $UR6
$5 = 0x0

Curiously enough, when the crash happens, the driver also reports a page fault:

NVRM: Xid (PCI:0000:04:00): 31, pid=105716, name=sm90_gemm_f32_f, Ch 00000008, intr 00000000. MMU Fault: ENGINE GRAPHICS GPCCLIENT_GCC faulted @ 0x7f79_78fff000. Fault is of type FAULT_PDE ACCESS_TYPE_VIRT_READ

However, I'm having trouble connecting reported fault address 0x7f79_78fff000 to the instruction inputs.

As far as I can tell, the data in the registers passed to the instruction is valid. I can't tell if the tensormap data is sensible, as it's not documented. The pointer reported in the driver log also does not seem to match anything passed to the instruction.

Any suggestions on what may be the root cause for UTMALDG.3D.MULTICAST triggering such a failure?

[thakkarV]

Usually this happens when your TMA descriptor is invalid. Need more details to help you debug:

• How are you constructing the TMA desc? With what input arguments?
• Does this crash happen at the very first issue of TMA? If not, when?
• What are the human readable values that are input to the PTX instruction for the coordinates, mbarrier pointer, TMA descriptor pointer?
• How are you passing the TMA desc to the kernel?
• What are you env details? Toolkit version? CUTLASS version? Host compiler? CMake commands to set up the build?

Please give as many details as you can

[Artem-B]

How are you constructing the TMA desc? With what input arguments?

I do not know the details. The issue is observed on a handful of gemm/**/sm90* tests compiled with clang. This particular crash and register dump is from the very first test in test/unit/gemm/device/sm90_gemm_f32_f32_f32_tensor_op_f32.cu here:

cutlass/test/unit/gemm/device/sm90_gemm_f32_f32_f32_tensor_op_f32.cu

Line 96 in bf9da7b

TEST(SM90_Device_Gemm_f32t_f32t_f32n_tensor_op_gmma_f32, 64x128x32_1x1x1_pingpong) {

It appears to be a regression from cutlass v3.4.x I'm updating from. Those test sources compile/run with v3.4 w/o issues, though they may not run the same tests.

Does this crash happen at the very first issue of TMA? If not, when?

I believe so. As far as I can tell it's the first cp.bulk*/UTMALDG* instruction executed by a "producer" thread. The crash happens in multiple threads, but I can reliably reproduce the crash even when I only leave a single thread running and early-exiting all other threads.

What are the human readable values that are input to the PTX instruction for the coordinates, mbarrier pointer, TMA descriptor pointer?

I only have line-info debug, and the build has some optimizations enabled (ptxas 12.6 crashes or hangs on unoptimized build for me) so I only have the registers passed to the failing instruction. The values are all in the cuda-gdb printouts above, but I'll include a copy below for convenience. Based on the SASS in https://godbolt.org/z/zhMrbhaWr what I have is

UTMALDG.3D.MULTICAST [UR8], [UR14], UR4, desc[UR6]

• UR8 - dst in shared memory. Appears to be valid and contains 32 LSBs of the generic address of the shared memory pointer passed to the kernel.

(cuda-gdb) p/x $UR8
$1 = 0x500

• UR14/15 point to the descriptor. I've also dumped some of the memory it points to.

(cuda-gdb) p/x $UR14
$2 = 0xedfff5c0
(cuda-gdb) p/x $UR15
$3 = 0x7ffd
(cuda-gdb) x/32x 0x7ffdedfff5c0
0x7ffdedfff5c0: 0xb5e00000      0x00007ffd      0x000563a0      0x00000001
0x7ffdedfff5d0: 0x00000000      0x00000000      0x00000000      0x00000000
0x7ffdedfff5e0: 0x00000003      0x00000003      0x00000000      0x00000000
0x7ffdedfff5f0: 0x00000000      0x1f000000      0x0000001f      0x00000000
0x7ffdedfff600: 0x00001000      0x00000000      0x00000400      0x00000000
0x7ffdedfff610: 0x00000000      0x00000000      0x00000000      0x00000000
0x7ffdedfff620: 0x00000000      0x00000000      0x00000000      0x00000000
0x7ffdedfff630: 0x00000000      0x00000000      0x00000000      0x00000000

• UR4 - multicast_mask

(cuda-gdb) p/x $UR4
$4 = 0x30000

• desc[UR6] - cache_hint

(cuda-gdb) p/x $UR6
$5 = 0x0

We appear to crash before we get to touch the barrier in subsequent SASS instructions.

How are you passing the TMA desc to the kernel?

No idea. Cutlass magic in cutlass test. The address we pass to the instruction does appear to match what we've got from the driver API call that created tensormap, but I should double check that.

What are you env details? Toolkit version? CUTLASS version? Host compiler? CMake commands to set up the build?

That's the tricky part. It's Google's internal build, using clang as CUDA compiler (close to LLVM/clang source tree HEAD, so it's more recent than any official release). The build has been ported to bazel, so cmake settings are not applicable. GPU is H100 SXM. The issue is not GPU-specific. The issue is reliably reproducible on multiple different machines and GPUs.

For what it's worth, barring a few minor issues, cutlass tests build and run quite well in our setup. These crashing tests are the only major regression I see updating from 3.4 to 3.6.

At the moment I can't tell if the issue is in our build (e.g our build may not be defining something correctly), the compiler (H100 support in NVPTX back-end is relatively new), or in cutlass (there are some issues with undefined behavior and nvcc/clang portability), and I need to get a better understanding of what's going on with that particular instruction and its inputs that leads to the exception.

• which of the instruction inputs is invalid, whether it's the direct value, or some indirect data in the descriptor?
• perhaps there's something else, and the instruction just happens to be the one to be executed when exception happens. If so what is that and how can I find it? I don't think that's the case, as it reliably happens if I single-step into it, but...
• Once I know which bit of data causes the crash, I'll start digging into what's wrong with that data. Is it corrupted, uninitialized, calculated incorrectly, etc.
• then it's just a matter of figuring out whodunnit. If it's clang, or our build, I'll fix it.