Skip to content
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

CUDA op getrows fails for long sequences #11189

Open
wants to merge 1 commit into
base: master
Choose a base branch
from
Open

CUDA op getrows fails for long sequences #11189

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
19 changes: 17 additions & 2 deletions ggml/src/ggml-cuda/getrows.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -118,13 +118,28 @@ static void get_rows_cuda_float(const ggml_tensor * src0, const ggml_tensor * sr
const size_t s12 = nb12 / ggml_element_size(src1);
//const size_t s13 = nb13 / ggml_element_size(src1);

k_get_rows_float<<<block_nums, block_dims, 0, stream>>>(
src0_dd, src1_dd, dst_dd,
// CUDA grids typically cannot exceed 65535 in .y dimension
static const int64_t MAX_GRID_Y = 65535;
for (int64_t startY = 0; startY < ne10; startY += MAX_GRID_Y) {
Comment on lines +122 to +123
Copy link
Collaborator

Choose a reason for hiding this comment

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

This code is incorrect. The grid y dimension uses 16 bits and ranges from 0 to 65535 (inclusive). So the correct stride would be 65536. With this code two threads per grid write to the same address (though this should result in identical results). The correct way to fix this would be to modify the CUDA kernel and have it iterate with a stride of 65536 over the y dimension. This will also avoid issues with the number of nodes in a CUDA graph varying depending on input parameters.

int64_t chunkSizeY = std::min<int64_t>(MAX_GRID_Y, ne10 - startY);

// launch chunkSizeY blocks in the Y-dimension instead of the entire ne10
dim3 block_nums(block_num_x, chunkSizeY, ne11 * ne12);

// pointer offset: since the kernel calculates src1[i10*s10 + ...] and dst[i10*s1 + ...],
// we shift src1_dd/dst_dd by (startY * stride) so that when i10=0 in the kernel,
// it corresponds to row "startY" in the overall space.
const int32_t * src1_dd_offset = src1_dd + startY * s10;
float * dst_dd_offset = dst_dd + startY * s1;

k_get_rows_float<<<block_nums, block_dims, 0, stream>>>(
src0_dd, src1_dd_offset, dst_dd_offset,
ne00, /*ne01, ne02, ne03,*/
/*ne10, ne11,*/ ne12, /*ne13,*/
/* s0,*/ s1, s2, s3,
/* nb00,*/ nb01, nb02, nb03,
s10, s11, s12/*, s13*/);
}

GGML_UNUSED(dst);
}
Expand Down
Loading