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]

It's possible, but without understanding how it's broken, it's hard to tell what's to fix. Unfortunately the problem is hard to reduce -- the user facing code is simple, but generates a rather large kernel that will never be a 1:1 match to what nvcc produces. It would indeed be one of the ways to approach the issue, but with cutlass, it's like searching for a needle in a haystack by comparing two huge haystacks. Doable in theory, but somewhat impractical. Cutlass also tends to make source code reproducers irreducible -- there are too many interconnected parts and attempts to reduce the problem using tools like creduce break down quickly, usually after they eliminate unused parts of the library headers, but that still leaves effectively all used source code and does not reduce the generated PTX code much.

If it was a regular instruction, I could reason which input is wrong and I could trace back the source. Here, I did check the inputs -- they look normal (to me), and accessing them manually works. I'm missing something implicit or indirect that I'm not aware of. My hunch is that it has something to do with the descriptor, but it's undocumented and opaque.

On my list of things to try next:

• verify that the descriptor content allocated on the host is the same one we're passing to the instruction.
• verify that descriptor creation parameters are the same between nvcc and clang-generated code.
• find or create a toy example of using TMA and see if I can reproduce the issue using the same TMA parameters that crash cutlass kernel. Then I can play with the inputs to the descriptor creation or to the instruction itself and see which one fails in a similar way.
• Long shot -- do compare NVCC and cutlass-generated code. Who knows, perhaps something will stand up and catch my eye.

If you have other hints or suggestions on possible ways to make progress, I'd appreciate them.

[d-k-b]

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.

@Artem-B, just as a sanity check, have you compared the exact compilation command line differences for the example .cu files between a cmake/clang-device run and the bazel/clang-device run? And also perhaps an cmake/nvcc run? CUTLASS builds can generate a number of definitions that are passed to the compilation command which affect which features are enabled. Ensuring these are the same would be a good first start.

[thakkarV]

verify that the descriptor content allocated on the host is the same one we're passing to the instruction.

Yes, this is the first thing I would ensure -- that the TMA descriptor passed in are bitwise accurate between NVCC and clang builds. After that, I would make sure the inputs to the PTX TMA instruction are also bitwise accurate.

If both of these are true, the next thing to ensure would be how the TMA desc itself is passed in. We pass those in via a new CUDA feature called grid constants. This has effects on the generated PTX and changes the programming model of how those kernel parameters are passed in.

[Artem-B]

That's a good point. I'll double check in case there's something new in v3.6 build that may affect these builds. I've been keeping internal cutlass builds since before cutlass 1.x, so I have a reasonable confidence that the build is sound in principle.

[Artem-B]

Yes, this is the first thing I would ensure -- that the TMA descriptor passed in are bitwise accurate between NVCC and clang builds. After that, I would make sure the inputs to the PTX TMA instruction are also bitwise accurate.

Based on the headers, it's expected to be 128 bytes in size and aligned on 64-byte boundary. If I understand the docs correctly, it incorporated the data pointer, so at least that part will not be bitwise identical, but it should be easy to identify.

Are there alignment requirements on the data pointer itself? Misaligned data pointer may potentially be one of the sources for the failure.