Skip to content

Fix CUDA illegal memory access bug in monotonic_rnnt #14

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

Merged
SimBe195 merged 2 commits into from
Feb 3, 2026
Merged

Fix CUDA illegal memory access bug in monotonic_rnnt #14

Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
afc1f44
Change dtype of `bts` from int to int64_t in grad kernel
SimBe195 Feb 2, 2026
cc01525
Change to int64_t also for act access index
SimBe195 Feb 2, 2026
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 i6_native_ops/monotonic_rnnt/include/gpu_rnnt_kernel.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -244,7 +244,7 @@ __global__ void compute_grad_kernel(Tp *grads, const Tp *const acts, const Tp *c
const int *const S_max, const int *const T_max, const int *const min_allowed_s,
const int *const max_allowed_s, const int *const V, const int blank_) {
int v = static_cast<int>(threadIdx.x);
int bts = static_cast<int>(blockIdx.x); // b, t, s packed
int64_t bts = static_cast<int64_t>(blockIdx.x); // b, t, s packed

int b = 0;
while (b < *B - 1 && denom_start_indices[b + 1] <= bts) {
Expand All @@ -259,7 +259,7 @@ __global__ void compute_grad_kernel(Tp *grads, const Tp *const acts, const Tp *c
const int *min_allowed_s_b = min_allowed_s + b * *T_max;
const int *max_allowed_s_b = max_allowed_s + b * *T_max;

int ts = bts - denom_start_indices[b];
int64_t ts = bts - denom_start_indices[b];
int t = ts / (S_b + 1);
int s = ts % (S_b + 1);

Expand Down