NVIDIA / NVIDIA/cutlass

Two small example bugs: wgrad split-k dumps the wrong tensor as Computed; FMHA diff-checker launches with block/grid swapped

Open Beginner friendly
#3,525 0 comments 0 reactions 0 assignees View on GitHub

Nobody has claimed this yet.

CUTLASS C++
Dominant language
C++
Stars
10.5k
Forks
2.1k
Avg merge
3d 11h
Merged PRs (30d)
7

Description

Description

Two small example defects spotted while running the examples on an RTX 5060 Ti (sm_120a, CUDA 12.8.1):

1. wgrad split-k workspace dump labels the wrong tensor as "Computed"

examples/30_wgrad_split_k/30_wgrad_split_k.cu ~line 621:

output_workspace << "Computed = \n" << tensor_c.host_view() << std::endl;

tensor_c is the zero-filled bias/filter-C input; the kernel output is tensor_d. With --save-workspace, anyone debugging gets a zeros dump labeled as the computed result.

Fix: print tensor_d.host_view().

2. Block/grid swapped in the FMHA reference diff-checker launches

examples/77_blackwell_fmha/reference/reference_abs_error.hpp (~lines 161-163 and 257-259) and examples/88_hopper_fmha/reference/reference_abs_error.hpp (~line 106):

dim3 block(256, 1, 1);
dim3 grid(1024, 1, 1);
reference_abs_diff_kernel<<<block, grid>>>(...);

This launches 256 blocks x 1024 threads instead of the intended 1024 x 256. Harmless for correctness here (grid-stride loops plus blockDim-sized shared reductions), but it teaches a wrong launch idiom in two widely copied files.

Fix: <<<grid, block>>>.

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

Inspect the indicated lines in examples/30_wgrad_split_k/30_wgrad_split_k.cu and the reference_abs_error.hpp files under examples/77_blackwell_fmha and examples/88_hopper_fmha. Compare the named tensors and kernel launch arguments with the surrounding code, then run the affected examples, including the split-k workspace dump, to verify the reported output and launch behavior are corrected.

Written by the indexing model from the issue text.

Assessment

Tech stack
cpp
Domain
hpc
Issue type
Bug
Difficulty
2/5
Estimated time
1-3 hours
Activity status
Active
Clarity
Clearly specified
Newbie friendliness
88/100

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.