[Bug]: GEMM Example 1 crashing

[Parxd]

Example 1 for GEMM 01_gemm_global_reg crashes during kernel for default problem size:

• M = 1024
• N = 1024
• K = 2048
• TileM = 256
• TileN = 128
• TileK = 64

System:
OS: Ubuntu 24.04.2
Device: RTX 3070 Mobile (SM86)
Driver: 560.35.03
CUDA Version: 12.6

Output:

kThreads: 128
RegA: RowMajor(16, 16)
RegB: ColMajor(16, 8)
RegC: RowMajor(16, 16)
IteratorA: numel = 524288, ChunkShape = (256, 64), stripe count = (1, 32)
IteratorB: numel = 262144, ChunkShape = (64, 128), stripe count = (32, 1)
blocks: [4, 8]

terminate called after throwing an instance of 'thrust::THRUST_200400_860_NS::system::system_error'
  what():  trivial_device_copy D->H failed: cudaErrorIllegalAddress: an illegal memory access was encountered
Aborted (core dumped)

[KuangjuX]

Example 1 for GEMM 01_gemm_global_reg crashes during kernel for default problem size:

• M = 1024
• N = 1024
• K = 2048
• TileM = 256
• TileN = 128
• TileK = 64

System: OS: Ubuntu 24.04.2 Device: RTX 3070 Mobile (SM86) Driver: 560.35.03 CUDA Version: 12.6

Output:

kThreads: 128
RegA: RowMajor(16, 16)
RegB: ColMajor(16, 8)
RegC: RowMajor(16, 16)
IteratorA: numel = 524288, ChunkShape = (256, 64), stripe count = (1, 32)
IteratorB: numel = 262144, ChunkShape = (64, 128), stripe count = (32, 1)
blocks: [4, 8]

terminate called after throwing an instance of 'thrust::THRUST_200400_860_NS::system::system_error'
  what():  trivial_device_copy D->H failed: cudaErrorIllegalAddress: an illegal memory access was encountered
Aborted (core dumped)

Thank you for your interest in our project and for providing the detailed information regarding the GEMM example. We appreciate your efforts in testing and reporting issues.

Currently, we do not have access to a machine with an RTX 3070, which limits our ability to directly troubleshoot this issue. However, based on the output you provided, one potential area to investigate is the tile sizes. You can try changing TileM, TileN, and TileK to see if that allows the test to pass.