microsoft / microsoft/TileFusion

[Bug]: GEMM Example 1 crashing

Open
#109 1 comment 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

Dominant language
Cuda
Stars
118
Forks
7
PR merge metrics
No merged PRs in 30d

Description

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)

Contributor guide

No contributing guide indexed for this repository

First steps

  1. Read the whole issue, then the project's contributing guide.
  2. Comment on the issue to say you are picking it up — it saves two people doing the same work.
  3. Fork the repository and make your change on a branch.
  4. Open a pull request that references the issue number.

Research direction

Start with the 01_gemm_global_reg GEMM example and reproduce the crash using its default problem size on the reported CUDA setup. Trace the kernel execution and the failing trivial_device_copy path; done means the example completes without the reported cudaErrorIllegalAddress crash.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Bug
Difficulty
4/5
Estimated time
3-5 days
Activity status
Stale
Clarity
Mostly clear
Newbie friendliness
35/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.