Skip to content

Fix non-deterministic results in k_grouped_fp8_gemm_nt_contiguous #375

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
Functionhx wants to merge 1 commit into
base: main
Choose a base branch
from
+10 −7
Open

Fix non-deterministic results in k_grouped_fp8_gemm_nt_contiguous #375

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
17 changes: 10 additions & 7 deletions deep_gemm/include/deep_gemm/impls/sm90_fp8_gemm_1d1d.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -191,26 +191,29 @@ sm90_fp8_gemm_1d1d_impl(__nv_fp8_e4m3* gmem_a_ptr, __nv_fp8_e4m3* gmem_b_ptr,
if (kGemmType == GemmType::KGroupedContiguous and last_group_idx != scheduler.current_group_idx) {
last_group_idx = scheduler.current_group_idx;

@ds-review-bot ds-review-bot Jul 9, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 suggestion: Root-cause wording is slightly imprecise. tensormap.replace is not a cp.async.bulk operation and is not part of any commit group, so the description's claim that replace and TMA loads 'compete within the same commit group' is not technically accurate. The real hazard is that the #343 fix placed the drain AFTER the replace, so the replace could run while a prior TMA descriptor read was still in flight (a missing happens-before). The fix itself is correct; please tighten the commit-message/comment framing to avoid future confusion.

🤖 v3

// Directly update current tensor map
// Drain all pending TMA data operations before modifying tensor map descriptors.
// This prevents a race where tensormap.replace operations and in-flight TMA loads
// compete within the same commit group, which the PTX ISA allows to resolve in any order.
cute::tma_desc_commit_group();
cute::tma_desc_wait_group();
__syncwarp(1U << lane_idx);

@ds-review-bot ds-review-bot Jul 9, 2026

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

🔵 suggestion: __syncwarp(1U << lane_idx) uses a single-lane mask, which is intentional here since this block runs under elect_one_sync() (only one lane active) and the __syncwarp acts purely as a compiler-ordering barrier. This diff removed the previous explanatory comment ('Only used to prevent ptxas from moving ... Shouldn't be needed otherwise, since we only use one thread'). Consider restoring a one-line comment on the single-lane __syncwarp so future readers don't mistake the mask for a bug.

🤖 v3


// Update current tensor map descriptors
const uint64_t current_k_offset = scheduler.current_k_cumsum;
ptx::tensor_map_replace_global_addr_in_smem(smem_tensor_map_a, gmem_a_ptr + current_k_offset * shape_m);
ptx::tensor_map_replace_global_addr_in_smem(smem_tensor_map_b, gmem_b_ptr + current_k_offset * shape_n);
ptx::tensor_map_replace_global_inner_dim_stride_in_smem(smem_tensor_map_a, scheduler.current_shape_k, scheduler.current_shape_k);
ptx::tensor_map_replace_global_inner_dim_stride_in_smem(smem_tensor_map_b, scheduler.current_shape_k, scheduler.current_shape_k);

// Make sure tensormaps are not used by TMA before updating GMEM.
// Fence the descriptor modifications
cute::tma_desc_commit_group();
cute::tma_desc_wait_group();
// Only used to prevent `ptxas` from moving the following GMEM stores before `cute::tma_desc_wait_group()`.
// Shouldn't be needed otherwise, since we only use one thread.

__syncwarp(1U << lane_idx);

// Copy updated descriptors to global memory and make them visible to TMA
*(gmem_tensor_map_a) = *(smem_tensor_map_a);
*(gmem_tensor_map_b) = *(smem_tensor_map_b);
ptx::tensor_map_release_gpu();

// Immediately acquire current tensor map
ptx::tensor_map_acquire_gpu(gmem_tensor_map_a);
ptx::tensor_map_acquire_gpu(gmem_tensor_map_b);
}
Expand Down