[Refactor] Remove dead code quantization 2 (#52821)
Signed-off-by: yewentao256 <zhyanwentao@126.com> Co-authored-by: mergify[bot] <37929162+mergify[bot]@users.noreply.github.com>
Wentao Ye committed
Aug 25, 2026 at 23:28 UTC
a5e0045433236d3b3c717ac2876578b4ed20b281
14 files changed
-526
csrc/libtorch_stable/quantization/gptq/matrix_view.cuh
-34
@@ -27,12 +27,6 @@ class MatrixView_half {
27
__device__ __forceinline__ half item(int row, int column) const {
28
return data[row * width + column];
29
}
30
- __device__ __forceinline__ half2 item_half2(int row, int column) const {
31
- return ((half2*)data)[(row * width + column) / 2];
32
- }
33
- __device__ __forceinline__ half2 item_half2half2(int row, int column) const {
34
- return __half2half2(data[row * width + column]);
35
- }
30
__device__ __forceinline__ const half* item_ptr(int row, int column) const {
31
return &data[row * width + column];
32
}
@@ -83,9 +77,6 @@ class MatrixView_half_rw {
77
__device__ __forceinline__ half item(int row, int column) const {
78
return data[row * width + column];
79
}
86
- __device__ __forceinline__ half2 item_half2(int row, int column) const {
87
- return ((half2*)data)[(row * width + column) / 2];
88
- }
80
__device__ __forceinline__ half* item_ptr(int row, int column) {
81
return &data[row * width + column];
82
}
@@ -141,31 +132,6 @@ class MatrixView_q4_row {
132
}
133
};
134
144
-class MatrixView_q4_column {
145
- public:
146
- const uint32_t* data;
147
- const int height;
148
- const int width;
149
-
150
- __device__ __forceinline__ MatrixView_q4_column(const uint32_t* data,
151
- const int height,
152
- const int width)
153
- : data(data), height(height), width(width) {}
154
-
155
- __device__ __forceinline__ int item(int row, int column) const {
156
- int shift = (row & 0x07) * 4;
157
- return (data[row / 8 * width + column] >> shift) & 0x0f;
158
- }
159
-
160
- __device__ __forceinline__ uint32_t item_uint32_t(int row, int column) {
161
- return data[row / 8 * width + column];
162
- }
163
- __device__ __forceinline__ const uint32_t* item_uint32_ptr(int row,
164
- int column) {
165
- return &data[row / 8 * width + column];
166
- }
167
-};
168
-
135
class MatrixView_q2_row {
136
public:
137
const uint32_t* data;
csrc/libtorch_stable/quantization/gptq/q_gemm.cu
-63
@@ -52,15 +52,6 @@ __host__ __forceinline__ hipblasStatus_t __compat_hipblasHgemm(
52
#define rocblas_hgemm __compat_hipblasHgemm
53
#endif
54
55
-__forceinline__ __device__ half2 dot22_8(half2 (&dq)[4], const half* a_ptr,
56
- const half2 g_result) {
57
- half2 result = {};
58
- const half2* a2_ptr = (const half2*)a_ptr;
59
-#pragma unroll
60
- for (int i = 0; i < 4; i++) result = __hfma2(dq[i], *a2_ptr++, result);
61
- return __hadd2(result, g_result);
62
-}
63
-
55
__forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
56
half2 result = {};
57
const half2* a2_ptr = (const half2*)a_ptr;
@@ -69,36 +60,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr) {
60
return __half2float(__low2half(result)) + __half2float(__high2half(result));
61
}
62
72
-__forceinline__ __device__ half2 dot22_8(half2 (&dq)[4], const half* a_ptr,
73
- const half2 g_result,
74
- const half qs_h) {
75
- half2 result = {};
76
- const half2* a2_ptr = (const half2*)a_ptr;
77
-#pragma unroll
78
- for (int i = 0; i < 4; i++) result = __hfma2(dq[i], *a2_ptr++, result);
79
- return __hfma2(result, __halves2half2(qs_h, qs_h), g_result);
80
-}
81
-
82
-__forceinline__ __device__ half2 dot22_16(half2 (&dq)[8], const half* a_ptr,
83
- const half2 g_result,
84
- const half qs_h) {
85
- half2 result = {};
86
- const half2* a2_ptr = (const half2*)a_ptr;
87
-#pragma unroll
88
- for (int i = 0; i < 8; i++) result = __hfma2(dq[i], *a2_ptr++, result);
89
- return __hfma2(result, __halves2half2(qs_h, qs_h), g_result);
90
-}
91
-
92
-__forceinline__ __device__ half2 dot22_32(half2 (&dq)[16], const half* a_ptr,
93
- const half2 g_result,
94
- const half qs_h) {
95
- half2 result = {};
96
- const half2* a2_ptr = (const half2*)a_ptr;
97
-#pragma unroll
98
- for (int i = 0; i < 16; i += 1) result = __hfma2(dq[i], *a2_ptr++, result);
99
- return __hfma2(result, __halves2half2(qs_h, qs_h), g_result);
100
-}
101
-
63
__forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr,
64
const float g_result,
65
const float qs_f) {
@@ -111,30 +72,6 @@ __forceinline__ __device__ float dot22_8_f(half2 (&dq)[4], const half* a_ptr,
72
return fma(result_f, qs_f, g_result);
73
}
74
114
-__forceinline__ __device__ float dot22_16_f(half2 (&dq)[8], const half* a_ptr,
115
- const float g_result,
116
- const float qs_f) {
117
- half2 result = {};
118
- const half2* a2_ptr = (const half2*)a_ptr;
119
-#pragma unroll
120
- for (int i = 0; i < 8; i++) result = __hfma2(dq[i], *a2_ptr++, result);
121
- float result_f =
122
- __half2float(__low2half(result)) + __half2float(__high2half(result));
123
- return fma(result_f, qs_f, g_result);
124
-}
125
-
126
-__forceinline__ __device__ float dot22_32_f(half2 (&dq)[16], const half* a_ptr,
127
- const float g_result,
128
- const float qs_f) {
129
- half2 result = {};
130
- const half2* a2_ptr = (const half2*)a_ptr;
131
-#pragma unroll
132
- for (int i = 0; i < 16; i += 1) result = __hfma2(dq[i], *a2_ptr++, result);
133
- float result_f =
134
- __half2float(__low2half(result)) + __half2float(__high2half(result));
135
- return fma(result_f, qs_f, g_result);
136
-}
137
-
75
__forceinline__ __device__ half dot22_8_h(half2 (&dq)[4], const half* a_ptr,
76
const half g_result,
77
const half qs_h) {
csrc/libtorch_stable/quantization/gptq/qdq_4.cuh
-42
@@ -28,48 +28,6 @@ __forceinline__ __device__ void shuffle_4bit_8(uint32_t* q, int stride) {
28
q[0] = qb;
29
}
30
31
-__forceinline__ __device__ void dequant_4bit_8(const uint32_t q_0,
32
- half2 (&dq)[4], int stride,
33
- const uint32_t zero) {
34
- const uint32_t c0 = 0x64006400;
35
- const half y16_ = __float2half_rn(1.0f / 16.0f);
36
- const half2 y16 = __halves2half2(y16_, y16_);
37
- const half_uint16 z1_(0xe400 | zero); // half(-1024.0f - zero);
38
- const half z16_ = __hsub(__int2half_rn(-64), __int2half_rn(zero));
39
- const half2 z1 = __half2half2(z1_.as_half);
40
- const half2 z16 = __half2half2(z16_);
41
-
42
- uint32_t qa = q_0;
43
- half2_uint32 q0((qa & 0x000f000f) | c0); // half2(q[ 0], q[ 1]) + 1024
44
- half2_uint32 q1((qa & 0x00f000f0) | c0); // half2(q[ 2], q[ 3]) * 16 + 1024
45
- qa >>= 8;
46
- half2_uint32 q2((qa & 0x000f000f) | c0); // half2(q[ 4], q[ 5]) + 1024
47
- half2_uint32 q3((qa & 0x00f000f0) | c0); // half2(q[ 6], q[ 7]) * 16 + 1024
48
-
49
- dq[0] = __hadd2(q0.as_half2, z1);
50
- dq[1] = __hfma2(q1.as_half2, y16, z16);
51
- dq[2] = __hadd2(q2.as_half2, z1);
52
- dq[3] = __hfma2(q3.as_half2, y16, z16);
53
-}
54
-
55
-__forceinline__ __device__ void dequant_4bit_8_prep_zero_scale(
56
- const uint32_t zero, const half scale, half2 (&z1z16)[2],
57
- half2 (&y1y16)[2]) {
58
- half_uint16 z1(0xe400 | zero); // half(-1024.0f - zero);
59
- half z16 = __hsub(__int2half_rn(-64), __int2half_rn(zero));
60
-
61
- half2 scale2 = __half2half2(scale);
62
-
63
- z1z16[0] = __hmul2(scale2, __half2half2(z1.as_half));
64
- z1z16[1] = __hmul2(scale2, __half2half2(z16));
65
-
66
- const half y1 = __float2half_rn(1.0f);
67
- const half y16 = __float2half_rn(1.0f / 16.0f);
68
-
69
- y1y16[0] = __hmul2(scale2, __half2half2(y1));
70
- y1y16[1] = __hmul2(scale2, __half2half2(y16));
71
-}
72
-
31
__forceinline__ __device__ void dequant_4bit_8_prep_zero(const uint32_t zero,
32
half2 (&z1z16)[2],
33
half2 (&y1y16)[2]) {
csrc/libtorch_stable/quantization/gptq/qdq_util.cuh
-9
@@ -22,15 +22,6 @@ union half_uint16 {
22
__device__ half_uint16(half val) : as_half(val) {}
23
};
24
25
-// Max_scale premultiplied by 1/256
26
-
27
-__forceinline__ __device__ half dq_scale(const int qs, const half max_scale) {
28
- int qs_i = qs + 1;
29
- half qs_h = __int2half_rn(qs_i * qs_i);
30
- qs_h = __hmul(qs_h, max_scale);
31
- return qs_h;
32
-}
33
-
25
__forceinline__ __device__ half dq(const int q, const int qzero,
26
const half scale) {
27
return __hmul(__int2half_rn(q - qzero), scale);
csrc/libtorch_stable/quantization/gptq_allspark/allspark_utils.cuh
-22
@@ -306,28 +306,6 @@ __device__ __forceinline__ void cp_async(const uint32_t smem_addr,
306
#endif
307
}
308
309
-template <int SIZE_IN_BYTES>
310
-__device__ __forceinline__ void cp_async_ca(const uint32_t smem_addr,
311
- const void* gmem_ptr,
312
- const int src_in_bytes,
313
- bool guard) {
314
- static_assert(
315
- (SIZE_IN_BYTES == 4 || SIZE_IN_BYTES == 8 || SIZE_IN_BYTES == 16),
316
- "Size is not supported");
317
-#if __CUDACC_VER_MAJOR__ >= 11 && __CUDA_ARCH__ >= 800
318
- asm volatile(
319
- "{.reg.pred p;\n"
320
- " setp.ne.b32 p, %4, 0;\n"
321
- #if __CUDACC_VER_MINOR__ >= 4
322
- " @p cp.async.ca.shared.global.L2::256B [%0], [%1], %2, %3;}\n"
323
- #else
324
- " @p cp.async.ca.shared.global [%0], [%1], %2, %3;}\n"
325
- #endif
326
- ::"r"(smem_addr),
327
- "l"(gmem_ptr), "n"(SIZE_IN_BYTES), "r"(src_in_bytes), "r"((int)guard));
328
-#endif
329
-}
330
-
309
__device__ __forceinline__ void cp_async_commit_group() {
310
#if __CUDACC_VER_MAJOR__ >= 11 && __CUDA_ARCH__ >= 800
311
asm volatile("cp.async.commit_group;\n");
csrc/libtorch_stable/quantization/machete/machete_prepacked_layout.cuh
-8
@@ -236,14 +236,6 @@ struct PrepackedLayoutBTemplate {
236
auto tiled_A = zipped_divide(make_layout(shape_mkl), tile);
237
return tiled_A.compose(ppblock_TV_to_NK(), _);
238
}
239
-
240
- // (N, K, L) -> ((athrid, val), (BlocksN, BlocksK), L)
241
- template <class Shape_NKL>
242
- CUTE_HOST_DEVICE static auto NKL_to_TVbNbK(Shape_NKL shape_mkl) {
243
- auto TVbNbK_to_NKL_layout = TVbNbK_to_NKL(shape_mkl);
244
- return blocked_product(ppblock_NK_to_TV(),
245
- make_layout(shape<1>(TVbNbK_to_NKL_layout)));
246
- }
239
};
240
241
}; // namespace machete
csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm.cuh
-61
@@ -84,67 +84,6 @@ struct cutlass_3x_gemm {
84
struct GemmKernel : public KernelType {};
85
};
86
87
-template <typename ElementAB_, typename ElementD_,
88
- template <typename, typename, typename> typename Epilogue_,
89
- typename TileShape, typename ClusterShape, typename KernelSchedule,
90
- typename EpilogueSchedule>
91
-struct cutlass_3x_gemm_sm100 {
92
- using ElementAB = ElementAB_;
93
- using LayoutA = cutlass::layout::RowMajor;
94
- static constexpr int AlignmentA =
95
- 128 / cutlass::sizeof_bits<ElementAB>::value;
96
-
97
- using LayoutB = cutlass::layout::ColumnMajor;
98
- static constexpr int AlignmentB =
99
- 128 / cutlass::sizeof_bits<ElementAB>::value;
100
-
101
- using ElementC = void;
102
- using LayoutC = cutlass::layout::RowMajor;
103
- static constexpr int AlignmentC =
104
- 128 / cutlass::sizeof_bits<ElementD_>::value;
105
-
106
- using ElementD = ElementD_;
107
- using LayoutD = cutlass::layout::RowMajor;
108
- static constexpr int AlignmentD = AlignmentC;
109
-
110
- using ElementAcc =
111
- typename std::conditional<std::is_same_v<ElementAB, int8_t>, int32_t,
112
- float>::type;
113
- using Epilogue = Epilogue_<ElementAcc, ElementD, TileShape>;
114
-
115
- // MMA type
116
- using ElementAccumulator = float;
117
-
118
- // Epilogue types
119
- using ElementBias = cutlass::half_t;
120
- using ElementCompute = float;
121
- using ElementAux = ElementD;
122
- using LayoutAux = LayoutD;
123
- using ElementAmax = float;
124
-
125
- using EVTCompute = typename Epilogue::EVTCompute;
126
-
127
- using CollectiveEpilogue =
128
- typename cutlass::epilogue::collective::CollectiveBuilder<
129
- cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp, TileShape,
130
- ClusterShape, cutlass::epilogue::collective::EpilogueTileAuto,
131
- ElementAccumulator, ElementCompute, ElementC, LayoutC, AlignmentC,
132
- ElementD, LayoutD, AlignmentD, EpilogueSchedule,
133
- EVTCompute>::CollectiveOp;
134
-
135
- using CollectiveMainloop =
136
- typename cutlass::gemm::collective::CollectiveBuilder<
137
- cutlass::arch::Sm100, cutlass::arch::OpClassTensorOp, ElementAB,
138
- LayoutA, AlignmentA, ElementAB, LayoutB, AlignmentB,
139
- ElementAccumulator, TileShape, ClusterShape,
140
- cutlass::gemm::collective::StageCountAutoCarveout<static_cast<int>(
141
- sizeof(typename CollectiveEpilogue::SharedStorage))>,
142
- KernelSchedule>::CollectiveOp;
143
-
144
- using GemmKernel = enable_sm100_to_sm120<cutlass::gemm::kernel::GemmUniversal<
145
- Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>>;
146
-};
147
-
87
template <typename ElementAB_, typename ElementD_,
88
template <typename, typename, typename> typename Epilogue_,
89
typename TileShape, typename ClusterShape, typename KernelSchedule,
csrc/quantization/w8a8/fp8/nvidia/quant_utils.cuh
-260
@@ -40,266 +40,6 @@ vec_conversion<c10::Float8_e4m3fn, float>(
40
#endif
41
}
42
43
- #if 0 // Disable the following code to reduce the binary size.
44
-// fp8 -> half
45
-template <>
46
-__inline__ __device__ uint16_t vec_conversion<uint16_t, uint8_t>(
47
- const uint8_t &a, const __nv_fp8_interpretation_t fp8_type) {
48
- __half_raw res = __nv_cvt_fp8_to_halfraw(a, fp8_type);
49
- return res.x;
50
-}
51
-
52
-// fp8x2 -> half2
53
-template <>
54
-__inline__ __device__ uint32_t vec_conversion<uint32_t, uint16_t>(
55
- const uint16_t &a, const __nv_fp8_interpretation_t fp8_type) {
56
- union {
57
- uint16_t u16[2];
58
- uint32_t u32;
59
- } tmp;
60
- __half2_raw res = __nv_cvt_fp8x2_to_halfraw2(a, fp8_type);
61
- tmp.u16[0] = res.x;
62
- tmp.u16[1] = res.y;
63
- return tmp.u32;
64
-}
65
-
66
-// fp8x4 -> half2x2
67
-template <>
68
-__inline__ __device__ uint2 vec_conversion<uint2, uint32_t>(
69
- const uint32_t &a, const __nv_fp8_interpretation_t fp8_type) {
70
- union {
71
- uint2 u32x2;
72
- uint32_t u32[2];
73
- } tmp;
74
- tmp.u32[0] = vec_conversion<uint32_t, uint16_t>((uint16_t)a, fp8_type);
75
- tmp.u32[1] =
76
- vec_conversion<uint32_t, uint16_t>((uint16_t)(a >> 16U), fp8_type);
77
- return tmp.u32x2;
78
-}
79
-
80
-// fp8x8 -> half2x4
81
-template <>
82
-__inline__ __device__ uint4 vec_conversion<uint4, uint2>(
83
- const uint2 &a, const __nv_fp8_interpretation_t fp8_type) {
84
- union {
85
- uint4 u64x2;
86
- uint2 u64[2];
87
- } tmp;
88
- tmp.u64[0] = vec_conversion<uint2, uint32_t>(a.x, fp8_type);
89
- tmp.u64[1] = vec_conversion<uint2, uint32_t>(a.y, fp8_type);
90
- return tmp.u64x2;
91
-}
92
-
93
-// fp8 -> __nv_bfloat16
94
-template <>
95
-__inline__ __device__ __nv_bfloat16 vec_conversion<__nv_bfloat16, uint8_t>(
96
- const uint8_t &a, const __nv_fp8_interpretation_t fp8_type) {
97
- // Note there is no direct convert function from fp8 to bf16.
98
- // fp8 -> half
99
- __half_raw res = __nv_cvt_fp8_to_halfraw(a, fp8_type);
100
- // half -> float -> bf16
101
- float tmp = half_to_float(res.x);
102
- return __float2bfloat16(tmp);
103
-}
104
-
105
-// fp8x2 -> __nv_bfloat162
106
-template <>
107
-__inline__ __device__ __nv_bfloat162 vec_conversion<__nv_bfloat162, uint16_t>(
108
- const uint16_t &a, const __nv_fp8_interpretation_t fp8_type) {
109
- __nv_bfloat162 res;
110
- res.x = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)a, fp8_type);
111
- res.y = vec_conversion<__nv_bfloat16, uint8_t>((uint8_t)(a >> 8U), fp8_type);
112
- return res;
113
-}
114
-
115
-// fp8x4 -> bf16_4_t
116
-template <>
117
-__inline__ __device__ bf16_4_t vec_conversion<bf16_4_t, uint32_t>(
118
- const uint32_t &a, const __nv_fp8_interpretation_t fp8_type) {
119
- bf16_4_t res;
120
- res.x = vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)a, fp8_type);
121
- res.y =
122
- vec_conversion<__nv_bfloat162, uint16_t>((uint16_t)(a >> 16U), fp8_type);
123
- return res;
124
-}
125
-
126
-// fp8x8 -> bf16_8_t
127
-template <>
128
-__inline__ __device__ bf16_8_t vec_conversion<bf16_8_t, uint2>(
129
- const uint2 &a, const __nv_fp8_interpretation_t fp8_type) {
130
- bf16_4_t tmp1, tmp2;
131
- tmp1 = vec_conversion<bf16_4_t, uint32_t>(a.x, fp8_type);
132
- tmp2 = vec_conversion<bf16_4_t, uint32_t>(a.y, fp8_type);
133
- bf16_8_t res;
134
- res.x = tmp1.x;
135
- res.y = tmp1.y;
136
- res.z = tmp2.x;
137
- res.w = tmp2.y;
138
- return res;
139
-}
140
-
141
-// fp8 -> float
142
-template <>
143
-__inline__ __device__ float
144
-vec_conversion<float, uint8_t>(const uint8_t &a,
145
- const __nv_fp8_interpretation_t fp8_type) {
146
- // fp8 -> half
147
- uint16_t tmp = vec_conversion<uint16_t, uint8_t>(a, fp8_type);
148
- // half -> float
149
- return half_to_float(tmp);
150
-}
151
-
152
-// fp8x2 -> float2
153
-template <>
154
-__inline__ __device__ float2 vec_conversion<float2, uint16_t>(
155
- const uint16_t &a, const __nv_fp8_interpretation_t fp8_type) {
156
- // fp8x2 -> half2
157
- uint32_t tmp = vec_conversion<uint32_t, uint16_t>(a, fp8_type);
158
- // half2 -> float2
159
- return half2_to_float2(tmp);
160
-}
161
-
162
-// fp8x4 -> float4
163
-template <>
164
-__inline__ __device__ Float4_ vec_conversion<Float4_, uint32_t>(
165
- const uint32_t &a, const __nv_fp8_interpretation_t fp8_type) {
166
- Float4_ res;
167
- res.x = vec_conversion<float2, uint16_t>((uint16_t)a, fp8_type);
168
- res.y = vec_conversion<float2, uint16_t>((uint16_t)(a >> 16U), fp8_type);
169
- return res;
170
-}
171
-
172
-// fp8x8 -> float8
173
-template <>
174
-__inline__ __device__ Float8_ vec_conversion<Float8_, uint2>(
175
- const uint2 &a, const __nv_fp8_interpretation_t fp8_type) {
176
- Float4_ tmp1, tmp2;
177
- tmp1 = vec_conversion<Float4_, uint32_t>(a.x, fp8_type);
178
- tmp2 = vec_conversion<Float4_, uint32_t>(a.y, fp8_type);
179
- Float8_ res;
180
- res.x = tmp1.x;
181
- res.y = tmp1.y;
182
- res.z = tmp2.x;
183
- res.w = tmp2.y;
184
- return res;
185
-}
186
-
187
-// half -> fp8
188
-template <>
189
-__inline__ __device__ uint8_t vec_conversion<uint8_t, uint16_t>(
190
- const uint16_t &a, const __nv_fp8_interpretation_t fp8_type) {
191
- __half_raw tmp;
192
- tmp.x = a;
193
- __nv_fp8_storage_t res =
194
- __nv_cvt_halfraw_to_fp8(tmp, __NV_SATFINITE, fp8_type);
195
- return (uint8_t)res;
196
-}
197
-
198
-// bf16 -> fp8
199
-template <>
200
-__inline__ __device__ uint8_t vec_conversion<uint8_t, __nv_bfloat16>(
201
- const __nv_bfloat16 &a, const __nv_fp8_interpretation_t fp8_type) {
202
- #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 800
203
- assert(false);
204
- #else
205
- __nv_fp8_storage_t res = __nv_cvt_bfloat16raw_to_fp8(
206
- __nv_bfloat16_raw(a), __NV_SATFINITE, fp8_type);
207
- return (uint8_t)res;
208
- #endif
209
-}
210
-
211
-// float -> fp8
212
-template <>
213
-__inline__ __device__ uint8_t vec_conversion<uint8_t, float>(
214
- const float &a, const __nv_fp8_interpretation_t fp8_type) {
215
- __nv_fp8_storage_t res = __nv_cvt_float_to_fp8(a, __NV_SATFINITE, fp8_type);
216
- return (uint8_t)res;
217
-}
218
-
219
-// fp8x4 -> float4
220
-template <>
221
-__inline__ __device__ float4 vec_conversion<float4, uint32_t>(
222
- const uint32_t &a, const __nv_fp8_interpretation_t fp8_type) {
223
- Float4_ tmp = vec_conversion<Float4_, uint32_t>(a, fp8_type);
224
- float4 res = make_float4(tmp.x.x, tmp.x.y, tmp.y.x, tmp.y.y);
225
- return res;
226
-}
227
-
228
-template <>
229
-__inline__ __device__ uint32_t vec_conversion<uint32_t, float2>(
230
- const float2 &a, const __nv_fp8_interpretation_t fp8_type) {
231
- union {
232
- half2 float16;
233
- uint32_t uint32;
234
- };
235
-
236
- float16 = __float22half2_rn(a);
237
- return uint32;
238
-}
239
-
240
-template <>
241
-__inline__ __device__ uint2 vec_conversion<uint2, Float4_>(
242
- const Float4_ &a, const __nv_fp8_interpretation_t fp8_type) {
243
- uint2 b;
244
- float2 val;
245
- val.x = a.x.x;
246
- val.y = a.x.y;
247
- b.x = vec_conversion<uint32_t, float2>(val, fp8_type);
248
-
249
- val.x = a.y.x;
250
- val.y = a.y.y;
251
- b.y = vec_conversion<uint32_t, float2>(val, fp8_type);
252
-
253
- return b;
254
-}
255
-
256
-template <>
257
-__inline__ __device__ float4 vec_conversion<float4, Float4_>(
258
- const Float4_ &a, const __nv_fp8_interpretation_t fp8_type) {
259
- float4 b;
260
- b.x = a.x.x;
261
- b.y = a.x.y;
262
- b.z = a.y.x;
263
- b.w = a.y.y;
264
- return b;
265
-}
266
-
267
-template <>
268
-__inline__ __device__ uint4 vec_conversion<uint4, Float8_>(
269
- const Float8_ &a, const __nv_fp8_interpretation_t fp8_type) {
270
- uint4 b;
271
- b.x = vec_conversion<uint32_t, float2>(a.x, fp8_type);
272
- b.y = vec_conversion<uint32_t, float2>(a.y, fp8_type);
273
- b.z = vec_conversion<uint32_t, float2>(a.z, fp8_type);
274
- b.w = vec_conversion<uint32_t, float2>(a.w, fp8_type);
275
- return b;
276
-}
277
-
278
-template <>
279
-__inline__ __device__ __nv_bfloat162 vec_conversion<__nv_bfloat162, float2>(
280
- const float2 &a, const __nv_fp8_interpretation_t fp8_type) {
281
- __nv_bfloat162 b;
282
- from_float(b, a);
283
- return b;
284
-}
285
-
286
-template <>
287
-__inline__ __device__ bf16_4_t vec_conversion<bf16_4_t, Float4_>(
288
- const Float4_ &a, const __nv_fp8_interpretation_t fp8_type) {
289
- bf16_4_t b;
290
- from_float(b, a);
291
- return b;
292
-}
293
-
294
-template <>
295
-__inline__ __device__ bf16_8_t vec_conversion<bf16_8_t, Float8_>(
296
- const Float8_ &a, const __nv_fp8_interpretation_t fp8_type) {
297
- bf16_8_t b;
298
- from_float(b, a);
299
- return b;
300
-}
301
- #endif
302
-
43
/* Scaled and vectorized conversions, for data exchange between high and low
44
precision domains Convention of the scale in API, e.g: FP8_data =
45
Quantization( High_Precision_data / scale ) s.t. Quantize(HP / scale) => FP8
vllm/model_executor/layers/quantization/compressed_tensors/schemes/compressed_tensors_w8a8_fp8.py
-4
@@ -38,9 +38,6 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
38
kFp8StaticChannelSym,
39
kFp8StaticTensorSym,
40
)
41
-from vllm.model_executor.layers.quantization.utils.w8a8_utils import (
42
- cutlass_block_fp8_supported,
43
-)
41
42
__all__ = ["CompressedTensorsW8A8Fp8"]
43
@@ -67,7 +64,6 @@ class CompressedTensorsW8A8Fp8(CompressedTensorsScheme):
64
self.weight_block_size = self.weight_quant.block_structure
65
66
if self.weight_block_size is not None:
70
- self.cutlass_block_fp8_supported = cutlass_block_fp8_supported()
67
self.use_aiter_and_is_supported = rocm_aiter_ops.is_linear_fp8_enabled()
68
assert not self.is_static_input_scheme
69
self.act_q_group_shape = GroupShape(1, self.weight_block_size[0])
vllm/model_executor/layers/quantization/fp8.py
-2
@@ -67,7 +67,6 @@ from vllm.model_executor.layers.quantization.utils.quant_utils import (
67
kFp8StaticTensorSym,
68
)
69
from vllm.model_executor.layers.quantization.utils.w8a8_utils import (
70
- cutlass_block_fp8_supported,
70
cutlass_fp8_supported,
71
normalize_e4m3fn_to_e4m3fnuz,
72
)
@@ -252,7 +251,6 @@ class Fp8LinearMethod(LinearMethodBase):
251
def __init__(self, quant_config: Fp8Config):
252
self.quant_config = quant_config
253
self.is_scale_e8m0 = getattr(quant_config, "is_scale_e8m0", False)
255
- self.cutlass_block_fp8_supported = cutlass_block_fp8_supported()
254
self.out_dtype = torch.get_default_dtype()
255
self.input_dtype = get_current_vllm_config().model_config.dtype
256
vllm/model_executor/layers/quantization/inc/schemes/inc_scheme.py
-16
@@ -9,7 +9,6 @@ if TYPE_CHECKING:
9
10
from vllm.model_executor.layers.fused_moe import FusedMoEMethodBase
11
from vllm.model_executor.layers.linear import LinearMethodBase
12
- from vllm.model_executor.layers.quantization import QuantizationMethods
12
13
from ..config_parser import INCLayerConfig
14
from ..inc import INCConfig
@@ -22,7 +21,6 @@ class INCScheme(ABC):
21
- can_handle(): when does this scheme apply?
22
- get_linear_method(): required — how to quantize Linear layers
23
- get_moe_method(): optional — how to quantize MoE layers
25
- - get_kvcache_method(): optional — how to quantize KV cache
24
25
Schemes that don't support MoE/KVCache inherit the default raise.
26
"""
@@ -56,20 +54,6 @@ class INCScheme(ABC):
54
f"Layer config: {layer_config}"
55
)
56
59
- def get_kvcache_method(
60
- self,
61
- config: "INCConfig",
62
- layer: "torch.nn.Module",
63
- prefix: str,
64
- layer_config: "INCLayerConfig",
65
- ) -> "QuantizationMethods":
66
- """Optional. Override if this scheme supports KV cache quantization.
67
- Default raises NotImplementedError."""
68
- raise NotImplementedError(
69
- f"{type(self).__name__} does not support KV cache quantization. "
70
- f"Layer config: {layer_config}"
71
- )
72
-
57
58
class INCLinearScheme(ABC):
59
@classmethod
vllm/model_executor/layers/quantization/quark/quark_moe.py
-2
@@ -1064,8 +1064,6 @@ class QuarkOCP_MX_MoEMethod(QuarkMoEMethod):
1064
else:
1065
self.input_dtype = None
1066
1067
- self.fp4_dtype = getattr(torch, "float4_e2m1fn_x2", None)
1068
-
1067
self.ocp_mx_scheme = OCP_MX_Scheme.from_quant_dtype(
1068
self.input_dtype, self.weight_dtype
1069
)
vllm/model_executor/layers/quantization/quark/schemes/quark_ocp_mx.py
-2
@@ -60,8 +60,6 @@ class QuarkOCP_MX(QuarkScheme):
60
input_quant_spec: dict[str, Any] | None,
61
dynamic_mxfp4_quant: bool = False,
62
):
63
- self.weight_quant_spec = weight_quant_spec
64
- self.input_quant_spec = input_quant_spec
63
self.dynamic_mxfp4_quant = dynamic_mxfp4_quant
64
self.weight_dtype = weight_quant_spec["dtype"].replace("fp", "mxfp")
65
self.input_dtype: str | None = None
vllm/model_executor/layers/quantization/utils/w8a8_utils.py
-1
@@ -38,7 +38,6 @@ def cutlass_group_gemm_supported() -> bool:
38
return ops.cutlass_group_gemm_supported(capability)
39
40
41
-CUTLASS_FP8_SUPPORTED = cutlass_fp8_supported()
41
CUTLASS_BLOCK_FP8_SUPPORTED = cutlass_block_fp8_supported()
42
43