[GH-ISSUE #34] Masking blocks in contiguous grouped GEMM does not work #22

New issue

Closed

opened 2026-10-26 21:10:41 +00:00 by bsakaguchi · 1 comment

bsakaguchi commented

2026-10-26 21:10:41 +00:00

Owner

Originally created by @hmwildermuth on GitHub (Mar 4, 2025).
Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/34

The documentation for m_grouped_gemm_fp8_fp8_bf16_nt_contiguous states that passing a value of -1 in m_indices will skip that block of 128 entries for the calculation.

However, this does not seem to be the case - there does not seem to be any code that does this, and passing -1 in fact triggers a out of bound memory access in testing.

The only relevant code that I can find seems to be the following:

    template <bool kIgnoreGroupedForGroupedContiguous=true>
    __device__ __forceinline__ uint32_t get_global_idx(const uint32_t shape_dim, const uint32_t block_size,
                                                       const uint32_t& block_idx, const uint32_t& m_block_idx=0) {
        if constexpr (kGemmType == GemmType::Normal) {
            return block_idx * block_size;
        } else if (kGemmType == GemmType::GroupedContiguous) {
            auto offset = kIgnoreGroupedForGroupedContiguous ? 0 : __ldg(grouped_layout + m_block_idx * BLOCK_M);
            return offset * shape_dim + block_idx * block_size;
        } else if (kGemmType == GemmType::GroupedMasked) {
            return curr_group_idx * shape_dim + block_idx * block_size;
        }
    }

However, kIgnoreGroupedForGroupedContiguous seems to be false at the relevant call sites in practice.

Originally created by @hmwildermuth on GitHub (Mar 4, 2025). Original GitHub issue: https://github.com/deepseek-ai/DeepGEMM/issues/34 The documentation for `m_grouped_gemm_fp8_fp8_bf16_nt_contiguous` states that passing a value of -1 in `m_indices` will skip that block of 128 entries for the calculation. However, this does not seem to be the case - there does not seem to be any code that does this, and passing -1 in fact triggers a out of bound memory access in testing. The only relevant code that I can find seems to be the following: ```c template <bool kIgnoreGroupedForGroupedContiguous=true> __device__ __forceinline__ uint32_t get_global_idx(const uint32_t shape_dim, const uint32_t block_size, const uint32_t& block_idx, const uint32_t& m_block_idx=0) { if constexpr (kGemmType == GemmType::Normal) { return block_idx * block_size; } else if (kGemmType == GemmType::GroupedContiguous) { auto offset = kIgnoreGroupedForGroupedContiguous ? 0 : __ldg(grouped_layout + m_block_idx * BLOCK_M); return offset * shape_dim + block_idx * block_size; } else if (kGemmType == GemmType::GroupedMasked) { return curr_group_idx * shape_dim + block_idx * block_size; } } ``` However, `kIgnoreGroupedForGroupedContiguous` seems to be false at the relevant call sites in practice.

bsakaguchi closed this issue

2026-10-26 21:10:42 +00:00

bsakaguchi commented

2026-10-26 21:10:43 +00:00

Author

Owner

@soundOfDestiny commented on GitHub (Mar 4, 2025):

The documentation is incorrect. I just fixed it in main, commit: github.com/deepseek-ai/DeepGEMM@ded740f736

Thanks for feedback!

@soundOfDestiny commented on GitHub (Mar 4, 2025): The documentation is incorrect. I just fixed it in main, commit: https://github.com/deepseek-ai/DeepGEMM/commit/ded740f736570c000b7063d7774f7c18b5d1ecd6 Thanks for feedback!

No labels

pull-request

No milestone

No project

No assignees

1 participant

Notifications

Due date

The due date is invalid or out of range. Please use the format "yyyy-mm-dd".

No due date set.

Dependencies

No dependencies set.

Reference

deepseek-ai/DeepGEMM#22

No description provided.

Rows
Columns