From 0c88cd01392c1073c7049a97d6328c7bba9b3947 Mon Sep 17 00:00:00 2001 From: shixianc <49539556+shixianc@users.noreply.github.com> Date: Sun, 15 Jun 2025 19:44:31 -0700 Subject: [PATCH] Fix illegal memory address when skipping -1 m indices (#113) Co-authored-by: Shixian Cui --- deep_gemm/include/deep_gemm/scheduler.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/deep_gemm/include/deep_gemm/scheduler.cuh b/deep_gemm/include/deep_gemm/scheduler.cuh index 81bfeba..69ea216 100644 --- a/deep_gemm/include/deep_gemm/scheduler.cuh +++ b/deep_gemm/include/deep_gemm/scheduler.cuh @@ -116,7 +116,7 @@ struct Scheduler { if constexpr (kGemmType == GemmType::Normal) { return block_idx * block_size; } else if constexpr (kGemmType == GemmType::GroupedContiguous) { - auto offset = kIgnoreGroupedForGroupedContiguous ? 0 : __ldg(grouped_layout + m_block_idx * BLOCK_M); + auto offset = kIgnoreGroupedForGroupedContiguous ? 0 : max(0, __ldg(grouped_layout + m_block_idx * BLOCK_M)); return offset * shape_dim + block_idx * block_size; } else if constexpr (kGemmType == GemmType::GroupedMasked) { return curr_group_idx * shape_dim + block_idx * block_size;