@@ -175,45 +175,45 @@ void percentileClipping_g16(half* g, float* gnorm_vec, int step, const int n) {
175175}
176176
177177void quantizeBlockwise_fp16 (float * code, half* A, float * absmax, unsigned char * out, int blocksize, const int n) {
178- quantizeBlockwise<half, 0 , General8bit>(code, A, absmax, out, NULL , 0 , blocksize, n);
178+ quantizeBlockwise<half, 0 , General8bit>(code, A, absmax, out, nullptr , 0 , blocksize, n);
179179}
180180
181181void quantizeBlockwise_fp16_fp4 (float * code, half* A, float * absmax, unsigned char * out, int blocksize, const int n) {
182- quantizeBlockwise<half, 0 , FP4>(NULL , A, absmax, out, NULL , 0 , blocksize, n);
182+ quantizeBlockwise<half, 0 , FP4>(nullptr , A, absmax, out, nullptr , 0 , blocksize, n);
183183}
184184
185185void quantizeBlockwise_fp16_nf4 (float * code, half* A, float * absmax, unsigned char * out, int blocksize, const int n) {
186- quantizeBlockwise<half, 0 , NF4>(NULL , A, absmax, out, NULL , 0 , blocksize, n);
186+ quantizeBlockwise<half, 0 , NF4>(nullptr , A, absmax, out, nullptr , 0 , blocksize, n);
187187}
188188
189189void quantizeBlockwise_bf16 (
190190 float * code, __nv_bfloat16* A, float * absmax, unsigned char * out, int blocksize, const int n
191191) {
192- quantizeBlockwise<__nv_bfloat16, 0 , General8bit>(code, A, absmax, out, NULL , 0 , blocksize, n);
192+ quantizeBlockwise<__nv_bfloat16, 0 , General8bit>(code, A, absmax, out, nullptr , 0 , blocksize, n);
193193}
194194
195195void quantizeBlockwise_bf16_fp4 (
196196 float * code, __nv_bfloat16* A, float * absmax, unsigned char * out, int blocksize, const int n
197197) {
198- quantizeBlockwise<__nv_bfloat16, 0 , FP4>(NULL , A, absmax, out, NULL , 0 , blocksize, n);
198+ quantizeBlockwise<__nv_bfloat16, 0 , FP4>(nullptr , A, absmax, out, nullptr , 0 , blocksize, n);
199199}
200200
201201void quantizeBlockwise_bf16_nf4 (
202202 float * code, __nv_bfloat16* A, float * absmax, unsigned char * out, int blocksize, const int n
203203) {
204- quantizeBlockwise<__nv_bfloat16, 0 , NF4>(NULL , A, absmax, out, NULL , 0 , blocksize, n);
204+ quantizeBlockwise<__nv_bfloat16, 0 , NF4>(nullptr , A, absmax, out, nullptr , 0 , blocksize, n);
205205}
206206
207207void quantizeBlockwise_fp32 (float * code, float * A, float * absmax, unsigned char * out, int blocksize, const int n) {
208- quantizeBlockwise<float , 0 , General8bit>(code, A, absmax, out, NULL , 0 , blocksize, n);
208+ quantizeBlockwise<float , 0 , General8bit>(code, A, absmax, out, nullptr , 0 , blocksize, n);
209209}
210210
211211void quantizeBlockwise_fp32_fp4 (float * code, float * A, float * absmax, unsigned char * out, int blocksize, const int n) {
212- quantizeBlockwise<float , 0 , FP4>(NULL , A, absmax, out, NULL , 0 , blocksize, n);
212+ quantizeBlockwise<float , 0 , FP4>(nullptr , A, absmax, out, nullptr , 0 , blocksize, n);
213213}
214214
215215void quantizeBlockwise_fp32_nf4 (float * code, float * A, float * absmax, unsigned char * out, int blocksize, const int n) {
216- quantizeBlockwise<float , 0 , NF4>(NULL , A, absmax, out, NULL , 0 , blocksize, n);
216+ quantizeBlockwise<float , 0 , NF4>(nullptr , A, absmax, out, nullptr , 0 , blocksize, n);
217217}
218218
219219void dequantizeBlockwise_fp16 (
@@ -225,13 +225,13 @@ void dequantizeBlockwise_fp16(
225225void dequantizeBlockwise_fp16_fp4 (
226226 float * code, unsigned char * A, float * absmax, half* out, int blocksize, const int n, cudaStream_t stream
227227) {
228- dequantizeBlockwise<half, FP4>(NULL , A, absmax, out, blocksize, n, stream);
228+ dequantizeBlockwise<half, FP4>(nullptr , A, absmax, out, blocksize, n, stream);
229229}
230230
231231void dequantizeBlockwise_fp16_nf4 (
232232 float * code, unsigned char * A, float * absmax, half* out, int blocksize, const int n, cudaStream_t stream
233233) {
234- dequantizeBlockwise<half, NF4>(NULL , A, absmax, out, blocksize, n, stream);
234+ dequantizeBlockwise<half, NF4>(nullptr , A, absmax, out, blocksize, n, stream);
235235}
236236
237237void dequantizeBlockwise_fp32 (
@@ -243,13 +243,13 @@ void dequantizeBlockwise_fp32(
243243void dequantizeBlockwise_fp32_fp4 (
244244 float * code, unsigned char * A, float * absmax, float * out, int blocksize, const int n, cudaStream_t stream
245245) {
246- dequantizeBlockwise<float , FP4>(NULL , A, absmax, out, blocksize, n, stream);
246+ dequantizeBlockwise<float , FP4>(nullptr , A, absmax, out, blocksize, n, stream);
247247}
248248
249249void dequantizeBlockwise_fp32_nf4 (
250250 float * code, unsigned char * A, float * absmax, float * out, int blocksize, const int n, cudaStream_t stream
251251) {
252- dequantizeBlockwise<float , NF4>(NULL , A, absmax, out, blocksize, n, stream);
252+ dequantizeBlockwise<float , NF4>(nullptr , A, absmax, out, blocksize, n, stream);
253253}
254254
255255void dequantizeBlockwise_bf16 (
@@ -261,13 +261,13 @@ void dequantizeBlockwise_bf16(
261261void dequantizeBlockwise_bf16_fp4 (
262262 float * code, unsigned char * A, float * absmax, __nv_bfloat16* out, int blocksize, const int n, cudaStream_t stream
263263) {
264- dequantizeBlockwise<__nv_bfloat16, FP4>(NULL , A, absmax, out, blocksize, n, stream);
264+ dequantizeBlockwise<__nv_bfloat16, FP4>(nullptr , A, absmax, out, blocksize, n, stream);
265265}
266266
267267void dequantizeBlockwise_bf16_nf4 (
268268 float * code, unsigned char * A, float * absmax, __nv_bfloat16* out, int blocksize, const int n, cudaStream_t stream
269269) {
270- dequantizeBlockwise<__nv_bfloat16, NF4>(NULL , A, absmax, out, blocksize, n, stream);
270+ dequantizeBlockwise<__nv_bfloat16, NF4>(nullptr , A, absmax, out, blocksize, n, stream);
271271}
272272
273273int igemmlt_32 (
@@ -323,13 +323,13 @@ void dequantizeBlockwise_fp16(
323323void dequantizeBlockwise_fp16_fp4 (
324324 float * code, unsigned char * A, float * absmax, sycl::half* out, int blocksize, const int n, sycl::queue* stream
325325) {
326- dequantizeBlockwise<sycl::half, FP4>(NULL , A, absmax, out, blocksize, n, stream);
326+ dequantizeBlockwise<sycl::half, FP4>(nullptr , A, absmax, out, blocksize, n, stream);
327327}
328328
329329void dequantizeBlockwise_fp16_nf4 (
330330 float * code, unsigned char * A, float * absmax, sycl::half* out, int blocksize, const int n, sycl::queue* stream
331331) {
332- dequantizeBlockwise<sycl::half, NF4>(NULL , A, absmax, out, blocksize, n, stream);
332+ dequantizeBlockwise<sycl::half, NF4>(nullptr , A, absmax, out, blocksize, n, stream);
333333}
334334
335335void dequantizeBlockwise_fp32 (
@@ -341,13 +341,13 @@ void dequantizeBlockwise_fp32(
341341void dequantizeBlockwise_fp32_fp4 (
342342 float * code, unsigned char * A, float * absmax, float * out, int blocksize, const int n, sycl::queue* stream
343343) {
344- dequantizeBlockwise<float , FP4>(NULL , A, absmax, out, blocksize, n, stream);
344+ dequantizeBlockwise<float , FP4>(nullptr , A, absmax, out, blocksize, n, stream);
345345}
346346
347347void dequantizeBlockwise_fp32_nf4 (
348348 float * code, unsigned char * A, float * absmax, float * out, int blocksize, const int n, sycl::queue* stream
349349) {
350- dequantizeBlockwise<float , NF4>(NULL , A, absmax, out, blocksize, n, stream);
350+ dequantizeBlockwise<float , NF4>(nullptr , A, absmax, out, blocksize, n, stream);
351351}
352352
353353void dequantizeBlockwise_bf16 (
@@ -361,14 +361,14 @@ void dequantizeBlockwise_bf16_fp4(
361361 float * code, unsigned char * A, float * absmax, sycl::ext::oneapi::bfloat16* out, int blocksize, const int n,
362362 sycl::queue* stream
363363) {
364- dequantizeBlockwise<sycl::ext::oneapi::bfloat16, FP4>(NULL , A, absmax, out, blocksize, n, stream);
364+ dequantizeBlockwise<sycl::ext::oneapi::bfloat16, FP4>(nullptr , A, absmax, out, blocksize, n, stream);
365365}
366366
367367void dequantizeBlockwise_bf16_nf4 (
368368 float * code, unsigned char * A, float * absmax, sycl::ext::oneapi::bfloat16* out, int blocksize, const int n,
369369 sycl::queue* stream
370370) {
371- dequantizeBlockwise<sycl::ext::oneapi::bfloat16, NF4>(NULL , A, absmax, out, blocksize, n, stream);
371+ dequantizeBlockwise<sycl::ext::oneapi::bfloat16, NF4>(nullptr , A, absmax, out, blocksize, n, stream);
372372}
373373
374374void gemv_4bit_inference_fp16 (
0 commit comments