[Refactor] Remove dead kernel code (#57621)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Wentao Ye committed
Sep 20, 2026 at 13:39 UTC
4868312128172a424a7cf4c90f25c53b49186346
12 files changed
+4
-559
csrc/cpu/sgl-kernels/common.h
-16
@@ -280,26 +280,10 @@ inline int get_thread_num() {
280
// balance payload across each thread
281
template <typename T>
282
inline void balance211(T n, T nth, T ith, T& n_start, T& n_end) {
283
-#if 0
284
- // onednn partition pattern
285
- T& n_my = n_end;
286
- if (nth <= 1 || n == 0) {
287
- n_start = 0;
288
- n_my = n;
289
- } else {
290
- T n1 = div_up(n, nth);
291
- T n2 = n1 - 1;
292
- T T1 = n - n2 * nth;
293
- n_my = ith < T1 ? n1 : n2;
294
- n_start = ith <= T1 ? ith*n1 : T1 * n1 + (ith - T1) * n2;
295
- }
296
- n_end += n_start;
297
-#else
283
// pytorch aten partition pattern
284
T n_my = div_up(n, nth);
285
n_start = ith * n_my;
286
n_end = std::min(n_start + n_my, n);
302
-#endif
287
}
288
289
template <typename func_t>
csrc/libtorch_stable/cuda_vec_utils.cuh
-4
@@ -320,10 +320,6 @@ __host__ __device__ __forceinline__ bool is_16byte_aligned(const void* ptr) {
320
return (reinterpret_cast<uintptr_t>(ptr) & 15) == 0;
321
}
322
323
-__host__ __device__ __forceinline__ bool is_32byte_aligned(const void* ptr) {
324
- return (reinterpret_cast<uintptr_t>(ptr) & 31) == 0;
325
-}
326
-
323
// ============================================================
324
// Packed type conversion and arithmetic
325
// ============================================================
csrc/libtorch_stable/moe/moe_wna16_utils.h
-8
@@ -34,10 +34,6 @@ class ScalarType<half> {
34
static __host__ __device__ float2 inline num22float2(const half2 x) {
35
return __half22float2(x);
36
}
37
-
38
- static __host__ __device__ half2 inline float22num2(const float2 x) {
39
- return __float22half2_rn(x);
40
- }
37
};
38
39
template <>
@@ -71,10 +67,6 @@ class ScalarType<nv_bfloat16> {
67
static __host__ __device__ float2 inline num22float2(const nv_bfloat162 x) {
68
return __bfloat1622float2(x);
69
}
74
-
75
- static __host__ __device__ nv_bfloat162 inline float22num2(const float2 x) {
76
- return __float22bfloat162_rn(x);
77
- }
70
#endif
71
};
72
csrc/libtorch_stable/quantization/gptq/matrix_view.cuh
-33
@@ -74,18 +74,9 @@ class MatrixView_half_rw {
74
const int width)
75
: data(data), height(height), width(width) {}
76
77
- __device__ __forceinline__ half item(int row, int column) const {
78
- return data[row * width + column];
79
- }
77
__device__ __forceinline__ half* item_ptr(int row, int column) {
78
return &data[row * width + column];
79
}
83
- __device__ __forceinline__ void set(int row, int column, half value) {
84
- data[row * width + column] = value;
85
- }
86
- __device__ __forceinline__ void set_half2(int row, int column, half2 value) {
87
- ((half2*)data)[(row * width + column) / 2] = value;
88
- }
80
81
__device__ __forceinline__ void set4(int row, int column, half v0, half v1,
82
half v2, half v3) {
@@ -113,14 +104,6 @@ class MatrixView_q4_row {
104
return (data[row * width / 8 + column / 8] >> shift) & 0x0f;
105
}
106
116
- __device__ __forceinline__ void item2(int (&items)[2], int row,
117
- int column) const {
118
- int shift = (column & 0x07) * 4;
119
- uint32_t d = data[row * width / 8 + column / 8] >> shift;
120
- items[0] = d & 0x0f;
121
- items[1] = (d >> 4) & 0x0f;
122
- }
123
-
107
__device__ __forceinline__ void item4(int (&items)[4], int row,
108
int column) const {
109
int shift = (column & 0x07) * 4;
@@ -148,14 +131,6 @@ class MatrixView_q2_row {
131
return (data[row * width / 16 + column / 16] >> shift) & 0x03;
132
}
133
151
- __device__ __forceinline__ void item2(int (&items)[2], int row,
152
- int column) const {
153
- int shift = (column & 0x0f) * 2;
154
- uint32_t d = data[row * width / 16 + column / 16] >> shift;
155
- items[0] = d & 0x03;
156
- items[1] = (d >> 2) & 0x03;
157
- }
158
-
134
__device__ __forceinline__ void item4(int (&items)[4], int row,
135
int column) const {
136
int shift = (column & 0x0f) * 2;
@@ -237,14 +212,6 @@ class MatrixView_q8_row {
212
return (data[row * width / 4 + column / 4] >> shift) & 0xff;
213
}
214
240
- __device__ __forceinline__ void item2(int (&items)[2], int row,
241
- int column) const {
242
- int shift = (column & 0x03) * 8;
243
- uint32_t d = data[row * width / 4 + column / 4] >> shift;
244
- items[0] = d & 0xff;
245
- items[1] = (d >> 8) & 0xff;
246
- }
247
-
215
__device__ __forceinline__ void item4(int (&items)[4], int row,
216
int column) const {
217
int shift = (column & 0x03) * 2;
csrc/libtorch_stable/quantization/marlin/marlin.cuh
+1
-22
@@ -67,14 +67,6 @@ __device__ inline void cp_async2_ca_pred(void* smem_ptr, const void* glob_ptr,
67
}
68
}
69
70
-__device__ inline void cp_async4_ca_pred(void* smem_ptr, const void* glob_ptr,
71
- bool pred = true) {
72
- if (pred) {
73
- reinterpret_cast<int4*>(smem_ptr)[0] =
74
- reinterpret_cast<const int4*>(glob_ptr)[0];
75
- }
76
-}
77
-
70
__device__ inline void cp_async4_pred(void* smem_ptr, const void* glob_ptr,
71
bool pred = true) {
72
if (pred) {
@@ -121,19 +113,6 @@ __device__ inline void cp_async2_ca_pred(void* smem_ptr, const void* glob_ptr,
113
"r"(smem), "l"(glob_ptr), "n"(BYTES));
114
}
115
124
-__device__ inline void cp_async4_ca_pred(void* smem_ptr, const void* glob_ptr,
125
- bool pred = true) {
126
- const int BYTES = 16;
127
- uint32_t smem = static_cast<uint32_t>(__cvta_generic_to_shared(smem_ptr));
128
- asm volatile(
129
- "{\n"
130
- " .reg .pred p;\n"
131
- " setp.ne.b32 p, %0, 0;\n"
132
- " @p cp.async.ca.shared.global [%1], [%2], %3;\n"
133
- "}\n" ::"r"((int)pred),
134
- "r"(smem), "l"(glob_ptr), "n"(BYTES));
135
-}
136
-
116
__device__ inline void cp_async4_pred(void* smem_ptr, const void* glob_ptr,
117
bool pred = true) {
118
const int BYTES = 16;
@@ -170,4 +149,4 @@ __device__ inline void cp_async_wait() {
149
150
} // namespace MARLIN_NAMESPACE_NAME
151
173
-#endif
\ No newline at end of file
152
+#endif
csrc/libtorch_stable/sampler.cu
-5
@@ -45,11 +45,6 @@ __global__ void apply_repetition_penalties_kernel(
45
}
46
}
47
48
-__device__ __forceinline__ auto convert_to_uint32(float x) -> uint32_t {
49
- uint32_t bits = __float_as_uint(x);
50
- return (bits & 0x80000000) ? bits : ~bits & 0x7fffffff;
51
-}
52
-
48
template <int step>
49
static inline __device__ uint32_t extractBinIdx(float x) {
50
if constexpr (step == 0) {
csrc/quantization/w8a8/fp8/amd/quant_utils.cuh
+1
-241
@@ -67,104 +67,6 @@ using fp8_type = __hip_fp8_e4m3_fnuz;
67
using fp8x2_type = __hip_fp8x2_e4m3_fnuz;
68
#endif
69
70
-// fp8 -> half
71
-template <>
72
-__inline__ __device__ uint16_t
73
-vec_conversion<uint16_t, uint8_t>(const uint8_t& a) {
74
- return __hip_cvt_fp8_to_halfraw(a, fp8_type::__default_interpret).x;
75
-}
76
-
77
-// fp8x2 -> half2
78
-template <>
79
-__inline__ __device__ uint32_t
80
-vec_conversion<uint32_t, uint16_t>(const uint16_t& a) {
81
- union {
82
- __half2_raw h2r;
83
- uint32_t ui32;
84
- } tmp;
85
- tmp.h2r = __hip_cvt_fp8x2_to_halfraw2(a, fp8_type::__default_interpret);
86
- return tmp.ui32;
87
-}
88
-
89
-// fp8x4 -> half2x2
90
-template <>
91
-__inline__ __device__ uint2 vec_conversion<uint2, uint32_t>(const uint32_t& a) {
92
- union {
93
- uint2 u32x2;
94
- uint32_t u32[2];
95
- } tmp;
96
- tmp.u32[0] = vec_conversion<uint32_t, uint16_t>((uint16_t)a);
97
- tmp.u32[1] = vec_conversion<uint32_t, uint16_t>((uint16_t)(a >> 16U));
98
- return tmp.u32x2;
99
-}
100
-
101
-// fp8x8 -> half2x4
102
-template <>
103
-__inline__ __device__ uint4 vec_conversion<uint4, uint2>(const uint2& a) {
104
- union {
105
- uint4 u64x2;
106
- uint2 u64[2];
107
- } tmp;
108
- tmp.u64[0] = vec_conversion<uint2, uint32_t>(a.x);
109
- tmp.u64[1] = vec_conversion<uint2, uint32_t>(a.y);
110
- return tmp.u64x2;
111
-}
112
-
113
-using __nv_bfloat16 = __hip_bfloat16;
114
-
115
-// fp8 -> __nv_bfloat16
116
-template <>
117
-__inline__ __device__ __nv_bfloat16
118
-vec_conversion<__nv_bfloat16, uint8_t>(const uint8_t& a) {
119
- fp8_type f8;
120
- f8.__x = a;
121
- return __float2bfloat16(static_cast<float>(f8));
122
-}
123
-
124
-using __nv_bfloat162 = __hip_bfloat162;
125
-
126
-// fp8x2 -> __nv_bfloat162
127
-template <>
128
-__inline__ __device__ __nv_bfloat162
129
-vec_conversion<__nv_bfloat162, uint16_t>(const uint16_t& a) {
130
- __nv_bfloat162 res;
131
- res.x = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)a);
132
- res.y = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)(a >> 8U));
133
- return res;
134
-}
135
-
136
-// fp8x4 -> bf16_4_t
137
-template <>
138
-__inline__ __device__ bf16_4_t
139
-vec_conversion<bf16_4_t, uint32_t>(const uint32_t& a) {
140
- bf16_4_t res;
141
- res.x = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)a);
142
- res.y = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)(a >> 16U));
143
- return res;
144
-}
145
-
146
-// fp8x8 -> bf16_8_t
147
-template <>
148
-__inline__ __device__ bf16_8_t vec_conversion<bf16_8_t, uint2>(const uint2& a) {
149
- bf16_4_t tmp1, tmp2;
150
- tmp1 = vec_conversion<bf16_4_t, uint32_t>(a.x);
151
- tmp2 = vec_conversion<bf16_4_t, uint32_t>(a.y);
152
- bf16_8_t res;
153
- res.x = tmp1.x;
154
- res.y = tmp1.y;
155
- res.z = tmp2.x;
156
- res.w = tmp2.y;
157
- return res;
158
-}
159
-
160
-// fp8 -> float
161
-template <>
162
-__inline__ __device__ float vec_conversion<float, uint8_t>(const uint8_t& a) {
163
- fp8_type f8;
164
- f8.__x = a;
165
- return static_cast<float>(f8);
166
-}
167
-
70
// fp8x2 -> float2
71
template <>
72
__inline__ __device__ float2
@@ -193,138 +95,6 @@ vec_conversion<float4, uint32_t>(const uint32_t& a) {
95
return res;
96
}
97
196
-// fp8x8 -> float8
197
-template <>
198
-__inline__ __device__ Float8_ vec_conversion<Float8_, uint2>(const uint2& a) {
199
- Float4_ tmp1, tmp2;
200
- tmp1 = vec_conversion<Float4_, uint32_t>(a.x);
201
- tmp2 = vec_conversion<Float4_, uint32_t>(a.y);
202
- Float8_ res;
203
- res.x = tmp1.x;
204
- res.y = tmp1.y;
205
- res.z = tmp2.x;
206
- res.w = tmp2.y;
207
- return res;
208
-}
209
-
210
-// half -> fp8
211
-template <>
212
-__inline__ __device__ uint8_t
213
-vec_conversion<uint8_t, uint16_t>(const uint16_t& a) {
214
- __half_raw tmp;
215
- tmp.x = a;
216
- return __hip_cvt_halfraw_to_fp8(tmp, fp8_type::__default_saturation,
217
- fp8_type::__default_interpret);
218
-}
219
-
220
-template <>
221
-__inline__ __device__ uint16_t
222
-vec_conversion<uint16_t, uint32_t>(const uint32_t& a) {
223
- union {
224
- uint32_t ui32;
225
- __half2_raw h2r;
226
- } tmp;
227
- tmp.ui32 = a;
228
- return __hip_cvt_halfraw2_to_fp8x2(tmp.h2r, fp8_type::__default_saturation,
229
- fp8_type::__default_interpret);
230
-}
231
-
232
-// bf16 -> fp8
233
-template <>
234
-__inline__ __device__ uint8_t
235
-vec_conversion<uint8_t, __nv_bfloat16>(const __nv_bfloat16& a) {
236
- return __hip_cvt_float_to_fp8(__bfloat162float(a),
237
- fp8_type::__default_saturation,
238
- fp8_type::__default_interpret);
239
-}
240
-
241
-// float -> fp8
242
-template <>
243
-__inline__ __device__ uint8_t vec_conversion<uint8_t, float>(const float& a) {
244
- return __hip_cvt_float_to_fp8(a, fp8_type::__default_saturation,
245
- fp8_type::__default_interpret);
246
-}
247
-
248
-// float2 -> half2
249
-template <>
250
-__inline__ __device__ uint32_t
251
-vec_conversion<uint32_t, float2>(const float2& a) {
252
- union {
253
- half2 float16;
254
- uint32_t uint32;
255
- };
256
-
257
- float16 = __float22half2_rn(a);
258
- return uint32;
259
-}
260
-
261
-// Float4 -> half2x2
262
-template <>
263
-__inline__ __device__ uint2 vec_conversion<uint2, Float4_>(const Float4_& a) {
264
- uint2 b;
265
- float2 val;
266
- val.x = a.x.x;
267
- val.y = a.x.y;
268
- b.x = vec_conversion<uint32_t, float2>(val);
269
-
270
- val.x = a.y.x;
271
- val.y = a.y.y;
272
- b.y = vec_conversion<uint32_t, float2>(val);
273
- return b;
274
-}
275
-
276
-// Float4 -> float4
277
-template <>
278
-__inline__ __device__ float4 vec_conversion<float4, Float4_>(const Float4_& a) {
279
- float4 b;
280
- b.x = a.x.x;
281
- b.y = a.x.y;
282
- b.z = a.y.x;
283
- b.w = a.y.y;
284
- return b;
285
-}
286
-
287
-// Float8 -> half2x4
288
-template <>
289
-__inline__ __device__ uint4 vec_conversion<uint4, Float8_>(const Float8_& a) {
290
- uint4 b;
291
- b.x = vec_conversion<uint32_t, float2>(a.x);
292
- b.y = vec_conversion<uint32_t, float2>(a.y);
293
- b.z = vec_conversion<uint32_t, float2>(a.z);
294
- b.w = vec_conversion<uint32_t, float2>(a.w);
295
- return b;
296
-}
297
-
298
-// float2 -> bfloat162
299
-template <>
300
-__inline__ __device__ __nv_bfloat162
301
-vec_conversion<__nv_bfloat162, float2>(const float2& a) {
302
- __nv_bfloat162 b = __float22bfloat162_rn(a);
303
- return b;
304
-}
305
-
306
-// Float4 -> bfloat162x2
307
-template <>
308
-__inline__ __device__ bf16_4_t
309
-vec_conversion<bf16_4_t, Float4_>(const Float4_& a) {
310
- bf16_4_t b;
311
- b.x = __float22bfloat162_rn(a.x);
312
- b.y = __float22bfloat162_rn(a.y);
313
- return b;
314
-}
315
-
316
-// Float8 -> bfloat162x4
317
-template <>
318
-__inline__ __device__ bf16_8_t
319
-vec_conversion<bf16_8_t, Float8_>(const Float8_& a) {
320
- bf16_8_t b;
321
- b.x = __float22bfloat162_rn(a.x);
322
- b.y = __float22bfloat162_rn(a.y);
323
- b.z = __float22bfloat162_rn(a.z);
324
- b.w = __float22bfloat162_rn(a.w);
325
- return b;
326
-}
327
-
98
/* Scaled and vectorized conversions, for data exchange between high and low
99
precision domains
100
@@ -335,6 +105,7 @@ vec_conversion<bf16_8_t, Float8_>(const Float8_& a) {
105
*/
106
107
using __nv_bfloat16 = __hip_bfloat16;
108
+using __nv_bfloat162 = __hip_bfloat162;
109
110
// fp8 -> __nv_bfloat16
111
template <>
@@ -613,17 +384,6 @@ scaled_vec_conversion<uint32_t, float4>(const float4& a, float scale) {
384
}
385
#endif // ENABLE_FP8
386
616
-template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
617
-__inline__ __device__ Tout convert(const Tin& x) {
618
- #ifdef ENABLE_FP8
619
- if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E4M3) {
620
- return vec_conversion<Tout, Tin>(x);
621
- }
622
- #endif
623
- assert(false);
624
- return {}; // Squash missing return statement warning
625
-}
626
-
387
template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
388
__inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) {
389
#ifdef ENABLE_FP8
csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh
-13
@@ -282,19 +282,6 @@ __inline__ __device__ float4 scaled_vec_conversion<float4, uint32_t>(
282
}
283
#endif // ENABLE_FP8
284
285
-template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
286
-__inline__ __device__ Tout convert(const Tin& x) {
287
- #if 0 // Disable the following code to reduce the binary size.
288
- if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E4M3) {
289
- return vec_conversion<Tout, Tin>(x, __NV_E4M3);
290
- } else if constexpr (kv_dt == Fp8KVCacheDataType::kFp8E5M2) {
291
- return vec_conversion<Tout, Tin>(x, __NV_E5M2);
292
- }
293
- #endif
294
- assert(false);
295
- __builtin_unreachable(); // Suppress missing return statement warning
296
-}
297
-
285
template <typename Tout, typename Tin, Fp8KVCacheDataType kv_dt>
286
__inline__ __device__ Tout scaled_convert(const Tin& x, const float scale) {
287
#ifdef ENABLE_FP8
csrc/rocm/attention.cu
-33
@@ -342,22 +342,6 @@ typedef union u64_cvt {
342
int64_t i64;
343
} _T8x8;
344
345
-__device__ __forceinline__ _B8x8 convert_b16x8(const _B16x8& input,
346
- _T8x8& Mtemp) {
347
- _T8x8 Qtmp8x8;
348
-
349
- for (int i = 0; i < 2; i++) {
350
- floatx4 q_out = {0, 0, 0, 0};
351
- q_out = gcn_mfma16x16x16_instr<_Float16, 0, 0, 0>(Mtemp.b64, input.xy[i],
352
- q_out);
353
- Qtmp8x8.b16x4[i * 2] =
354
- __builtin_amdgcn_cvt_pk_fp8_f32(q_out[0], q_out[1], 0, false);
355
- Qtmp8x8.b16x4[i * 2 + 1] =
356
- __builtin_amdgcn_cvt_pk_fp8_f32(q_out[2], q_out[3], 0, false);
357
- }
358
- return Qtmp8x8.b8x8;
359
-}
360
-
345
__device__ float warpReduceMax(float val) {
346
for (int offset = warpSize / 2; offset > 0; offset /= 2) {
347
val = max(
@@ -2583,23 +2567,6 @@ __device__ __forceinline__ float to_float(const T& inp) {
2567
}
2568
}
2569
2586
-template <typename T>
2587
-__device__ __forceinline__ float to_float_b16(const bit16_t& inp) {
2588
- union tmpcvt {
2589
- bit16_t u;
2590
- _Float16 f;
2591
- __hip_bfloat16 b;
2592
- } t16;
2593
- t16.u = inp;
2594
- if constexpr (std::is_same<T, _Float16>::value) {
2595
- return (float)t16.f;
2596
- } else if constexpr (std::is_same<T, __hip_bfloat16>::value) {
2597
- return __bfloat162float(t16.b);
2598
- } else {
2599
- static_assert(false, "unsupported 16b dtype");
2600
- }
2601
-}
2602
-
2570
template <typename T>
2571
__device__ __forceinline__ T from_float(const float& inp) {
2572
if constexpr (std::is_same<T, _Float16>::value) {
csrc/rocm/moe_q_gemm_rdna3.cu
-15
@@ -67,21 +67,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
67
return result;
68
}
69
70
-__forceinline__ __device__ float dot22_8_f(float (&dq)[8],
71
- const bf16_t* a_ptr) {
72
- float result = 0.0f;
73
- #pragma unroll
74
- for (int i = 0; i < 4; i++) {
75
- uint32_t aw;
76
- __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
77
- float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
78
- float a_y = __uint_as_float(aw & 0xFFFF0000u);
79
- result = __fmaf_rn(dq[2 * i + 0], a_x, result);
80
- result = __fmaf_rn(dq[2 * i + 1], a_y, result);
81
- }
82
- return result;
83
-}
84
-
70
__forceinline__ __device__ void atomic_add_pk4_f16(half* addr, half2 v01,
71
half2 v23) {
72
unsigned long long* addr_u = reinterpret_cast<unsigned long long*>(addr);
csrc/rocm/q_gemm_rdna3.cu
+2
-58
@@ -115,61 +115,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
115
return result;
116
}
117
118
-__forceinline__ __device__ float dot22_8_f(bf162_t (&dq)[4],
119
- const bf16_t* a_ptr) {
120
- // RDNA3 (gfx1100) lacks a packed bf16 FMA: there is no v_pk_fma_bf16 in
121
- // the gfx11 ISA (it only landed on CDNA3+ / gfx94x and later). hipcc
122
- // therefore lowers __hfma2(bf162_t, bf162_t, bf162_t) to a serialised
123
- // fallback (single-element FMAs or fp32 round-trips), which empirically
124
- // runs ~2× the cycle count of v_pk_fma_f16 on the same VALU. The bf16
125
- // decode path was paying that tax in full, scaling linearly with M (the
126
- // fp16 path scales sub-linearly because its v_pk_fma_f16 is full rate
127
- // and the kernel becomes memory-bound).
128
- //
129
- // Fix: widen bf16 → fp32 explicitly (a left-shift by 16, free in VGPRs)
130
- // and accumulate with v_fma_f32, which IS full rate on RDNA3. Same FMA
131
- // count, but each FMA is fast. Bonus: the accumulator is now fp32
132
- // throughout instead of bf16, which is also numerically more accurate
133
- // (no compounding bf16-rounding inside the dot loop).
134
- float result = 0.0f;
135
- #pragma unroll
136
- for (int i = 0; i < 4; i++) {
137
- uint32_t aw, dw;
138
- __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
139
- __builtin_memcpy(&dw, &dq[i], sizeof(uint32_t));
140
- // bf16 in low 16 bits → fp32 by left-shifting into the upper half.
141
- // bf16 in high 16 bits → already aligned with fp32's upper half.
142
- float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
143
- float a_y = __uint_as_float(aw & 0xFFFF0000u);
144
- float d_x = __uint_as_float((dw & 0xFFFFu) << 16);
145
- float d_y = __uint_as_float(dw & 0xFFFF0000u);
146
- result = __fmaf_rn(d_x, a_x, result);
147
- result = __fmaf_rn(d_y, a_y, result);
148
- }
149
- return result;
150
-}
151
-
152
-// fp32-input dot product: paired with dequant_4bit_8_bf16_f32 which already
153
-// produces fp32 dq[8]. Saves the bf16→fp32 widening that the bf162_t
154
-// overload above does for dq (still need to widen A from bf16). Wins more
155
-// at high N: the bf162_t version's per-call widening cost scales with the
156
-// number of dequants × M_COUNT × 4 dot calls; the fp32 version pays only
157
-// for A widening (M_COUNT × 4 × 4 widens, half as many).
158
-__forceinline__ __device__ float dot22_8_f(float (&dq)[8],
159
- const bf16_t* a_ptr) {
160
- float result = 0.0f;
161
- #pragma unroll
162
- for (int i = 0; i < 4; i++) {
163
- uint32_t aw;
164
- __builtin_memcpy(&aw, a_ptr + 2 * i, sizeof(uint32_t));
165
- float a_x = __uint_as_float((aw & 0xFFFFu) << 16);
166
- float a_y = __uint_as_float(aw & 0xFFFF0000u);
167
- result = __fmaf_rn(dq[2 * i + 0], a_x, result);
168
- result = __fmaf_rn(dq[2 * i + 1], a_y, result);
169
- }
170
- return result;
171
-}
172
-
118
// ---------------------------------------------------------------------------
119
// Packed atomic-add via CAS-loop on a 64-bit word (4 fp16/bf16 lanes per CAS).
120
// RDNA3 (gfx11) does NOT have native v_global_atomic_pk_add_f16 / _bf16 (those
@@ -327,9 +272,8 @@ __global__ void gemm_q4_kernel_rdna3(const T* __restrict__ a,
272
273
// Per-column dequant constants. We hold one set of (z, y) pairs per column.
274
// fp16 uses the exllama (z1z16, y1y16) double-pair to enable the upper-
330
- // nibble-*16 trick. bf16 uses fp32 scalars (z, y) because the dequant
331
- // produces fp32 directly — see prep_zero_scale_bf16_f32 / the FMA
332
- // bypass for the missing v_pk_fma_bf16 on gfx11.
275
+ // nibble-*16 trick. bf16 uses fp32 scalars (z, y) for the inline v_dot2
276
+ // bias correction.
277
half2 z1z16_h[4][2], y1y16_h[4][2];
278
float z_b_f[4], y_b_f[4];
279
csrc/rocm/qdq_4_rdna3.cuh
-111
@@ -113,63 +113,6 @@ __forceinline__ __device__ void dequant_4bit_8_fp16(uint32_t qa, half2 (&dq)[4],
113
// bf16 path
114
// ---------------------------------------------------------------------------
115
116
-// Bit-trick magic for bf16:
117
-// bf16(128) == 0x4300 (sign 0, exp 134, mantissa 0).
118
-// For nibble n in [0..15], bits [3:0] of mantissa hold n exactly because
119
-// bf16's ULP at 128 is 1 (mantissa step = 2^(7-7) = 1). So
120
-// ((qa & 0x000F000F) | 0x43004300) bitcasts to bfloat162(128+n_lo, 128+n_hi).
121
-//
122
-// Because bf16's mantissa is only 7 bits, we cannot use the fp16 "upper nibble
123
-// * 16" trick. Instead each pair of nibbles is shifted down to [3:0]/[19:16]
124
-// via a single 4/8/12-bit right-shift before the OR. That costs one extra
125
-// shift per pair vs fp16, but keeps the FMA structure identical.
126
-__forceinline__ __device__ void prep_zero_scale_bf16(uint32_t zero,
127
- bf16_t scale,
128
- bf162_t& z_prep,
129
- bf162_t& y_prep) {
130
- // z = scale * -(128 + zero); y = scale.
131
- float scale_f = __bfloat162float(scale);
132
- float zf = -(128.0f + (float)zero) * scale_f;
133
- bf16_t zb = __float2bfloat16(zf);
134
- z_prep = __bfloat162bfloat162(zb);
135
- y_prep = __bfloat162bfloat162(scale);
136
-}
137
-
138
-__forceinline__ __device__ void dequant_4bit_8_bf16(uint32_t qa,
139
- bf162_t (&dq)[4],
140
- bf162_t z_prep,
141
- bf162_t y_prep) {
142
- const uint32_t c0 = 0x43004300;
143
-
144
- union {
145
- uint32_t u;
146
- bf162_t b2;
147
- } q0, q1, q2, q3;
148
- q0.u = ((qa >> 0) & 0x000F000F) | c0; // bf162(128+q[0], 128+q[1])
149
- q1.u = ((qa >> 4) & 0x000F000F) | c0; // bf162(128+q[2], 128+q[3])
150
- q2.u = ((qa >> 8) & 0x000F000F) | c0; // bf162(128+q[4], 128+q[5])
151
- q3.u = ((qa >> 12) & 0x000F000F) | c0; // bf162(128+q[6], 128+q[7])
152
-
153
- // dq = q_b * scale + (-(128+zero)*scale) = (q - zero) * scale
154
- dq[0] = __hfma2(q0.b2, y_prep, z_prep);
155
- dq[1] = __hfma2(q1.b2, y_prep, z_prep);
156
- dq[2] = __hfma2(q2.b2, y_prep, z_prep);
157
- dq[3] = __hfma2(q3.b2, y_prep, z_prep);
158
-}
159
-
160
-// ---------------------------------------------------------------------------
161
-// bf16-input → fp32-output dequant (RDNA3 scalar path).
162
-//
163
-// RDNA3 (gfx1100) has no v_pk_fma_bf16; packed bf16 FMA lowers to a slow
164
-// fallback. Rather than computing dq in bf16 and widening at FMA time in
165
-// the dot product, we widen to fp32 here once (a free left-shift by 16) and
166
-// emit the (q - zero) * scale FMA directly in fp32. This:
167
-// * Replaces 4× slow bf16 packed FMA with 8× fast fp32 FMA per int32.
168
-// * Eliminates 4× bf16→fp32 widens that the dot product would do.
169
-// * Keeps the dot product accumulator in fp32 without a roundtrip.
170
-//
171
-// Output: fp32 dq[8], one element per K position (consumed by the
172
-// fp32-overload of dot22_8_f in q_gemm_rdna3.cu).
116
__forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero,
117
bf16_t scale,
118
float& z_prep,
@@ -179,60 +122,6 @@ __forceinline__ __device__ void prep_zero_scale_bf16_f32(uint32_t zero,
122
y_prep = scale_f;
123
}
124
182
-// Pure-q dequant for the M_COUNT=1 factored path: outputs the unscaled fp32
183
-// values 128+nibble, without folding scale/zero. The caller folds scale/zb
184
-// into the accumulator outside the inner loop using a precomputed sum_a,
185
-// which saves ~27% of the FMA count vs the per-col-dequant approach above
186
-// (only beneficial at M_COUNT=1; break-even at M_COUNT=2).
187
-//
188
-// Cost: 0 FMAs (pure bit-trick + as_float reinterprets).
189
-__forceinline__ __device__ void dequant_4bit_8_bf16_q_only(uint32_t qa,
190
- float (&q_f32)[8]) {
191
- const uint32_t c0 = 0x43004300;
192
- const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0;
193
- const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0;
194
- const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0;
195
- const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0;
196
- q_f32[0] = __uint_as_float((q0 & 0xFFFFu) << 16);
197
- q_f32[1] = __uint_as_float(q0 & 0xFFFF0000u);
198
- q_f32[2] = __uint_as_float((q1 & 0xFFFFu) << 16);
199
- q_f32[3] = __uint_as_float(q1 & 0xFFFF0000u);
200
- q_f32[4] = __uint_as_float((q2 & 0xFFFFu) << 16);
201
- q_f32[5] = __uint_as_float(q2 & 0xFFFF0000u);
202
- q_f32[6] = __uint_as_float((q3 & 0xFFFFu) << 16);
203
- q_f32[7] = __uint_as_float(q3 & 0xFFFF0000u);
204
-}
205
-
206
-__forceinline__ __device__ void dequant_4bit_8_bf16_f32(uint32_t qa,
207
- float (&dq)[8],
208
- float z_prep,
209
- float y_prep) {
210
- const uint32_t c0 = 0x43004300;
211
- const uint32_t q0 = ((qa >> 0) & 0x000F000F) | c0;
212
- const uint32_t q1 = ((qa >> 4) & 0x000F000F) | c0;
213
- const uint32_t q2 = ((qa >> 8) & 0x000F000F) | c0;
214
- const uint32_t q3 = ((qa >> 12) & 0x000F000F) | c0;
215
- // bf16(128+nibble) bits → fp32(128+nibble) bits via left-shift by 16
216
- // (just zero-extends the mantissa from 7 to 23 bits; exponent preserved).
217
- const float q0x = __uint_as_float((q0 & 0xFFFFu) << 16);
218
- const float q0y = __uint_as_float(q0 & 0xFFFF0000u);
219
- const float q1x = __uint_as_float((q1 & 0xFFFFu) << 16);
220
- const float q1y = __uint_as_float(q1 & 0xFFFF0000u);
221
- const float q2x = __uint_as_float((q2 & 0xFFFFu) << 16);
222
- const float q2y = __uint_as_float(q2 & 0xFFFF0000u);
223
- const float q3x = __uint_as_float((q3 & 0xFFFFu) << 16);
224
- const float q3y = __uint_as_float(q3 & 0xFFFF0000u);
225
- // dq[i] = q_f32 * scale + (-(128+zero)*scale) = (nibble - zero) * scale
226
- dq[0] = __fmaf_rn(q0x, y_prep, z_prep);
227
- dq[1] = __fmaf_rn(q0y, y_prep, z_prep);
228
- dq[2] = __fmaf_rn(q1x, y_prep, z_prep);
229
- dq[3] = __fmaf_rn(q1y, y_prep, z_prep);
230
- dq[4] = __fmaf_rn(q2x, y_prep, z_prep);
231
- dq[5] = __fmaf_rn(q2y, y_prep, z_prep);
232
- dq[6] = __fmaf_rn(q3x, y_prep, z_prep);
233
- dq[7] = __fmaf_rn(q3y, y_prep, z_prep);
234
-}
235
-
125
} // namespace gptq_rdna3
126
} // namespace vllm
127