diff --git a/csrc/activation_kernels.cu b/csrc/activation_kernels.cu index 303433392c32..d8e9fa894777 100644 --- a/csrc/activation_kernels.cu +++ b/csrc/activation_kernels.cu @@ -79,9 +79,16 @@ __global__ void act_and_mul_kernel( scalar_t* __restrict__ out, // [..., d] const scalar_t* __restrict__ input, // [..., 2, d] const int d, const float limit) { - const scalar_t* x_ptr = input + blockIdx.x * 2 * d; + // `blockIdx.x` is `unsigned int`; multiplying by `d` (`int`) keeps the + // result in 32 bits, which overflows once `blockIdx.x * 2 * d` exceeds + // INT_MAX (about 2.15 billion). For large hidden sizes this corrupts the + // pointer arithmetic and reads/writes the wrong memory. Promote the index + // to int64_t first; matches the pattern already used in + // `swigluoai_and_mul_kernel` below. See issue #42860. + const int64_t token_idx = blockIdx.x; + const scalar_t* x_ptr = input + token_idx * 2 * d; const scalar_t* y_ptr = x_ptr + d; - scalar_t* out_ptr = out + blockIdx.x * d; + scalar_t* out_ptr = out + token_idx * d; if constexpr (use_vec) { using cuda_t = typename CUDATypeConverter::Type; @@ -313,9 +320,12 @@ template