Skip to content

[ROCm][GPTQ][Bugfix] Fix GPTQ GEMM kernel output zeroing race condition #30719

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
vllm-bot merged 10 commits into from
Dec 29, 2025
Merged

[ROCm][GPTQ][Bugfix] Fix GPTQ GEMM kernel output zeroing race condition #30719

Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
7326fef
fix output zeroing race condition in GPTQ GEMM kernels
AndreasKaratzas Dec 15, 2025
a36191a
fix output zeroing race condition in GPTQ GEMM kernels
AndreasKaratzas Dec 15, 2025
036f6b8
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 15, 2025
0cceed1
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 18, 2025
275a131
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 24, 2025
829de12
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 24, 2025
bab137c
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 26, 2025
657698c
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 27, 2025
82b6f71
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 28, 2025
6003ed8
Merge remote-tracking branch 'origin/main' into fix/gptq-rocm
AndreasKaratzas Dec 28, 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
28 changes: 1 addition & 27 deletions csrc/quantization/gptq/q_gemm.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -233,11 +233,6 @@ __global__ void gemm_half_q_half_gptq_4bit_kernel(
// Zero output
if (n >= size_n) return;

if (blockIdx.z == 0) {
for (int m = 0; m < m_count; m++)
*((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
}

__syncthreads();

// Find initial group
Expand Down Expand Up @@ -372,11 +367,6 @@ __global__ void gemm_half_q_half_gptq_2bit_kernel(
// Zero output
if (n >= size_n) return;

if (blockIdx.z == 0) {
for (int m = 0; m < m_count; m++)
*((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
}

__syncthreads();

// Find initial group
Expand Down Expand Up @@ -494,11 +484,6 @@ __global__ void gemm_half_q_half_gptq_3bit_kernel(
// Zero output
if (n >= size_n) return;

if (blockIdx.z == 0) {
for (int m = 0; m < m_count; m++)
*((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
}

__syncthreads();

// Find initial group
Expand Down Expand Up @@ -623,11 +608,6 @@ __global__ void gemm_half_q_half_gptq_8bit_kernel(
// Zero output
if (n >= size_n) return;

if (blockIdx.z == 0) {
for (int m = 0; m < m_count; m++)
*((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
}

__syncthreads();

// Find initial group
Expand Down Expand Up @@ -1224,9 +1204,6 @@ __global__ void gemm_half_q_half_alt_4bit_kernel(
__halves2half2(__int2half_rn(val & 0xF), __int2half_rn(val >> 4));
}

if (blockIdx.z == 0) {
for (int m = 0; m < b_end; m++) mul[(b + m) * width + w] = __int2half_rn(0);
}
__syncthreads();

int i = width * h + w;
Expand Down Expand Up @@ -1319,9 +1296,6 @@ __global__ void gemm_half_q_half_alt_8bit_kernel(
}
}

if (blockIdx.z == 0) {
for (int m = 0; m < b_end; m++) mul[(b + m) * width + w] = __int2half_rn(0);
}
__syncthreads();

int i = width * h + w;
Expand Down Expand Up @@ -1857,7 +1831,7 @@ torch::Tensor gptq_gemm(torch::Tensor a, torch::Tensor b_q_weight,
bool use_exllama, bool use_v2_format, int64_t bit) {
const at::cuda::OptionalCUDAGuard device_guard(device_of(a));
auto options = torch::TensorOptions().dtype(a.dtype()).device(a.device());
at::Tensor c = torch::empty({a.size(0), b_q_weight.size(1)}, options);
at::Tensor c = torch::zeros({a.size(0), b_q_weight.size(1)}, options);
at::Tensor temp_dq = torch::empty(
{b_q_weight.size(0) * 32 / bit, b_q_weight.size(1)}, options);

Expand Down