Skip to content

Add grid-stride + ROCm cap to grad_mean kernel #6078

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
q10 wants to merge 1 commit into
base: main
Choose a base branch
from
Open

Add grid-stride + ROCm cap to grad_mean kernel #6078

Show file tree
Hide file tree
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
14 changes: 13 additions & 1 deletion fbgemm_gpu/codegen/training/backward/embedding_backward_split_grad_template.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -158,14 +158,23 @@ __global__ __launch_bounds__(kMaxThreads) void grad_mean{{ vdesc }}_kernel(
{% endif %}
) {
int32_t T = D_offsets.size(0) - 1;
auto b_t = blockIdx.x * blockDim.y + threadIdx.y;
[[maybe_unused]] int32_t b;
int32_t t;
const auto total_B = offsets.size(0) - 1;

// On ROCm the launch caps the grid to stay within the HIP 2^32
// threads-per-launch limit, so we grid-stride to cover the full workload.
// On CUDA the grid is not capped and the loop body runs once per warp.
#ifdef USE_ROCM
for (auto b_t = blockIdx.x * blockDim.y + threadIdx.y;
b_t < total_B;
b_t += blockDim.y * gridDim.x) {
#else
auto b_t = blockIdx.x * blockDim.y + threadIdx.y;
if (b_t >= total_B) {
return;
}
#endif

{% if vbe %}
const auto info = reinterpret_cast<const uint32_t*>(&b_t_map[b_t])[0];
Expand Down Expand Up @@ -205,6 +214,9 @@ __global__ __launch_bounds__(kMaxThreads) void grad_mean{{ vdesc }}_kernel(
grad_out_vec.store(&shifted_grad_output_mean[d * 4]);
}
}
#ifdef USE_ROCM
} // for b_t (grid-stride loop, ROCm only)
#endif
}

////////////////////////////////////////////////////////////////////////////////
Expand Down
5 changes: 4 additions & 1 deletion fbgemm_gpu/codegen/training/backward/embedding_backward_split_template.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -952,7 +952,10 @@ Tensor {{ embedding_cuda_op }}(
const int grad_mean_warp_size = at::cuda::warp_size();
FBGEMM_LAUNCH_KERNEL(
(grad_mean{{ vdesc }}_kernel<grad_t, index_t>),
div_round_up(total_B, kMaxThreads / grad_mean_warp_size),
utils::cuda::cap_grid_dim_x(
div_round_up(total_B, kMaxThreads / grad_mean_warp_size),
kMaxThreads,
at::cuda::getCurrentCUDAStream()),
dim3(grad_mean_warp_size, kMaxThreads / grad_mean_warp_size),
0,
at::cuda::getCurrentCUDAStream(),
Expand Down
Loading