A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 10, 21:30 UTC
[ROCm][GPTQ][Bugfix] Fix GPTQ GEMM kernel output zeroing race condition (#30719)
Signed-off-by: Andreas Karatzas <akaratza@amd.com>
Andreas Karatzas committed
Dec 29, 2025 at 03:13 UTC
3ecfdc377681064f85684bb1ce5cc44c82e69b03
1 file changed
+1
-27
csrc/quantization/gptq/q_gemm.cu
+1
-27
@@ -233,11 +233,6 @@ __global__ void gemm_half_q_half_gptq_4bit_kernel(
233
// Zero output
234
if (n >= size_n) return;
235
236
- if (blockIdx.z == 0) {
237
- for (int m = 0; m < m_count; m++)
238
- *((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
239
- }
240
-
236
__syncthreads();
237
238
// Find initial group
@@ -372,11 +367,6 @@ __global__ void gemm_half_q_half_gptq_2bit_kernel(
367
// Zero output
368
if (n >= size_n) return;
369
375
- if (blockIdx.z == 0) {
376
- for (int m = 0; m < m_count; m++)
377
- *((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
378
- }
379
-
370
__syncthreads();
371
372
// Find initial group
@@ -494,11 +484,6 @@ __global__ void gemm_half_q_half_gptq_3bit_kernel(
484
// Zero output
485
if (n >= size_n) return;
486
497
- if (blockIdx.z == 0) {
498
- for (int m = 0; m < m_count; m++)
499
- *((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
500
- }
501
-
487
__syncthreads();
488
489
// Find initial group
@@ -623,11 +608,6 @@ __global__ void gemm_half_q_half_gptq_8bit_kernel(
608
// Zero output
609
if (n >= size_n) return;
610
626
- if (blockIdx.z == 0) {
627
- for (int m = 0; m < m_count; m++)
628
- *((uint64_t*)c_.item_ptr(offset_m + m, n)) = 0;
629
- }
630
-
611
__syncthreads();
612
613
// Find initial group
@@ -1224,9 +1204,6 @@ __global__ void gemm_half_q_half_alt_4bit_kernel(
1204
__halves2half2(__int2half_rn(val & 0xF), __int2half_rn(val >> 4));
1205
}
1206
1227
- if (blockIdx.z == 0) {
1228
- for (int m = 0; m < b_end; m++) mul[(b + m) * width + w] = __int2half_rn(0);
1229
- }
1207
__syncthreads();
1208
1209
int i = width * h + w;
@@ -1319,9 +1296,6 @@ __global__ void gemm_half_q_half_alt_8bit_kernel(
1296
}
1297
}
1298
1322
- if (blockIdx.z == 0) {
1323
- for (int m = 0; m < b_end; m++) mul[(b + m) * width + w] = __int2half_rn(0);
1324
- }
1299
__syncthreads();
1300
1301
int i = width * h + w;
@@ -1857,7 +1831,7 @@ torch::Tensor gptq_gemm(torch::Tensor a, torch::Tensor b_q_weight,
1831
bool use_exllama, bool use_v2_format, int64_t bit) {
1832
const at::cuda::OptionalCUDAGuard device_guard(device_of(a));
1833
auto options = torch::TensorOptions().dtype(a.dtype()).device(a.device());
1860
- at::Tensor c = torch::empty({a.size(0), b_q_weight.size(1)}, options);
1834
+ at::Tensor c = torch::zeros({a.size(0), b_q_weight.size(1)}, options);
1835
at::Tensor temp_dq = torch::empty(
1836
{b_q_weight.size(0) * 32 / bit, b_q_weight.size(1)}, options);
1837