Skip to content

Update gather_mm.cu #7889

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
FirstStep0 wants to merge 2 commits into
base: master
Choose a base branch
from
Open

Update gather_mm.cu #7889

Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
979351a
Update gather_mm.cu
FirstStep0 Jun 9, 2025
7321961
Update gather_mm.cu
FirstStep0 Jun 18, 2025
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
4 changes: 2 additions & 2 deletions src/array/cuda/gather_mm.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -113,7 +113,7 @@ __global__ void GatherMMScatterKernel(
for (unsigned int outloop = 0; outloop < out_len; outloop += 32) {
DType out_reg = static_cast<DType>(0.0f); // thread private
const unsigned int l = laneId;
if (l < out_len) {
if (l + outloop < out_len) {
// iterate over elements of a row of A
for (unsigned int i = 0; i < a_tile; i++) {
const DType a_val = sh_A[local_row * sh_a_tile + i];
Expand Down Expand Up @@ -170,7 +170,7 @@ __global__ void GatherMMScatterKernel2(
for (unsigned int outloop = 0; outloop < out_len; outloop += 32) {
DType out_reg = static_cast<DType>(0.0f); // thread private
const unsigned int l = laneId;
if (l < out_len) {
if (l + outloop < out_len) {
const DType b_val = B[row_b * out_len + (outloop + l)];
/* iterate over elements of a row of A */
for (unsigned int i = 0; i < a_tile; i++) {
Expand Down