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 15, 17:15 UTC
[Kernel] Fix conflicting macro names for gguf kernels (#15456)
Signed-off-by: SzymonOzog <szymon.ozog@gmail.com>
Szymon Ożóg committed
Mar 25, 2025 at 14:50 UTC
a60816002724925c8eee5978d75ebec2775e17e8
2 files changed
+90
-90
csrc/quantization/gguf/gguf_kernel.cu
+10
-10
@@ -375,25 +375,25 @@ torch::Tensor ggml_moe_a8(torch::Tensor X, // input
375
int64_t ggml_moe_get_block_size(int64_t type) {
376
switch (type) {
377
case 2:
378
- return MMQ_X_Q4_0;
378
+ return MOE_X_Q4_0;
379
case 3:
380
- return MMQ_X_Q4_1;
380
+ return MOE_X_Q4_1;
381
case 6:
382
- return MMQ_X_Q5_0;
382
+ return MOE_X_Q5_0;
383
case 7:
384
- return MMQ_X_Q5_1;
384
+ return MOE_X_Q5_1;
385
case 8:
386
- return MMQ_X_Q8_0;
386
+ return MOE_X_Q8_0;
387
case 10:
388
- return MMQ_X_Q2_K;
388
+ return MOE_X_Q2_K;
389
case 11:
390
- return MMQ_X_Q3_K;
390
+ return MOE_X_Q3_K;
391
case 12:
392
- return MMQ_X_Q4_K;
392
+ return MOE_X_Q4_K;
393
case 13:
394
- return MMQ_X_Q5_K;
394
+ return MOE_X_Q5_K;
395
case 14:
396
- return MMQ_X_Q6_K;
396
+ return MOE_X_Q6_K;
397
}
398
return 0;
399
}
csrc/quantization/gguf/moe.cuh
+80
-80
@@ -129,12 +129,12 @@ static __device__ __forceinline__ void moe_q(
129
}
130
131
#if defined(USE_ROCM)
132
- #define MMQ_X_Q4_0 64
133
- #define MMQ_Y_Q4_0 128
132
+ #define MOE_X_Q4_0 64
133
+ #define MOE_Y_Q4_0 128
134
#define NWARPS_Q4_0 8
135
#else
136
- #define MMQ_X_Q4_0 4
137
- #define MMQ_Y_Q4_0 32
136
+ #define MOE_X_Q4_0 4
137
+ #define MOE_Y_Q4_0 32
138
#define NWARPS_Q4_0 4
139
#endif
140
@@ -149,8 +149,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q4_0, 2)
149
const int exp_stride, const int ncols_x, const int nrows_x,
150
const int ncols_y, const int nrows_y, const int nrows_dst,
151
const int top_k) {
152
- const int mmq_x = MMQ_X_Q4_0;
153
- const int mmq_y = MMQ_Y_Q4_0;
152
+ const int mmq_x = MOE_X_Q4_0;
153
+ const int mmq_y = MOE_Y_Q4_0;
154
const int nwarps = NWARPS_Q4_0;
155
156
moe_q<scalar_t, QK4_0, QR4_0, QI4_0, true, block_q4_0, mmq_x, mmq_y, nwarps,
@@ -167,8 +167,8 @@ static void ggml_moe_q4_0_q8_1_cuda(
167
const int exp_stride, const int ncols_x, const int nrows_x,
168
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
169
const int tokens_post_padded, cudaStream_t stream) {
170
- int mmq_x = MMQ_X_Q4_0;
171
- int mmq_y = MMQ_Y_Q4_0;
170
+ int mmq_x = MOE_X_Q4_0;
171
+ int mmq_y = MOE_Y_Q4_0;
172
int nwarps = NWARPS_Q4_0;
173
174
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -190,12 +190,12 @@ static void ggml_moe_q4_0_q8_1_cuda(
190
}
191
192
#if defined(USE_ROCM)
193
- #define MMQ_X_Q4_1 64
194
- #define MMQ_Y_Q4_1 128
193
+ #define MOE_X_Q4_1 64
194
+ #define MOE_Y_Q4_1 128
195
#define NWARPS_Q4_1 8
196
#else
197
- #define MMQ_X_Q4_1 4
198
- #define MMQ_Y_Q4_1 32
197
+ #define MOE_X_Q4_1 4
198
+ #define MOE_Y_Q4_1 32
199
#define NWARPS_Q4_1 4
200
#endif
201
@@ -210,8 +210,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q4_1, 2)
210
const int exp_stride, const int ncols_x, const int nrows_x,
211
const int ncols_y, const int nrows_y, const int nrows_dst,
212
const int top_k) {
213
- const int mmq_x = MMQ_X_Q4_1;
214
- const int mmq_y = MMQ_Y_Q4_1;
213
+ const int mmq_x = MOE_X_Q4_1;
214
+ const int mmq_y = MOE_Y_Q4_1;
215
const int nwarps = NWARPS_Q4_1;
216
217
moe_q<scalar_t, QK4_1, QR4_1, QI4_1, true, block_q4_1, mmq_x, mmq_y, nwarps,
@@ -228,8 +228,8 @@ static void ggml_moe_q4_1_q8_1_cuda(
228
const int exp_stride, const int ncols_x, const int nrows_x,
229
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
230
const int tokens_post_padded, cudaStream_t stream) {
231
- int mmq_x = MMQ_X_Q4_1;
232
- int mmq_y = MMQ_Y_Q4_1;
231
+ int mmq_x = MOE_X_Q4_1;
232
+ int mmq_y = MOE_Y_Q4_1;
233
int nwarps = NWARPS_Q4_1;
234
235
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -251,12 +251,12 @@ static void ggml_moe_q4_1_q8_1_cuda(
251
}
252
253
#if defined(USE_ROCM)
254
- #define MMQ_X_Q5_0 64
255
- #define MMQ_Y_Q5_0 128
254
+ #define MOE_X_Q5_0 64
255
+ #define MOE_Y_Q5_0 128
256
#define NWARPS_Q5_0 8
257
#else
258
- #define MMQ_X_Q5_0 4
259
- #define MMQ_Y_Q5_0 32
258
+ #define MOE_X_Q5_0 4
259
+ #define MOE_Y_Q5_0 32
260
#define NWARPS_Q5_0 4
261
#endif
262
@@ -271,8 +271,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q5_0, 2)
271
const int exp_stride, const int ncols_x, const int nrows_x,
272
const int ncols_y, const int nrows_y, const int nrows_dst,
273
const int top_k) {
274
- const int mmq_x = MMQ_X_Q5_0;
275
- const int mmq_y = MMQ_Y_Q5_0;
274
+ const int mmq_x = MOE_X_Q5_0;
275
+ const int mmq_y = MOE_Y_Q5_0;
276
const int nwarps = NWARPS_Q5_0;
277
278
moe_q<scalar_t, QK5_0, QR5_0, QI5_0, false, block_q5_0, mmq_x, mmq_y, nwarps,
@@ -289,8 +289,8 @@ static void ggml_moe_q5_0_q8_1_cuda(
289
const int exp_stride, const int ncols_x, const int nrows_x,
290
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
291
const int tokens_post_padded, cudaStream_t stream) {
292
- const int mmq_x = MMQ_X_Q5_0;
293
- const int mmq_y = MMQ_Y_Q5_0;
292
+ const int mmq_x = MOE_X_Q5_0;
293
+ const int mmq_y = MOE_Y_Q5_0;
294
const int nwarps = NWARPS_Q5_0;
295
296
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -312,12 +312,12 @@ static void ggml_moe_q5_0_q8_1_cuda(
312
}
313
314
#if defined(USE_ROCM)
315
- #define MMQ_X_Q5_1 64
316
- #define MMQ_Y_Q5_1 128
315
+ #define MOE_X_Q5_1 64
316
+ #define MOE_Y_Q5_1 128
317
#define NWARPS_Q5_1 8
318
#else
319
- #define MMQ_X_Q5_1 4
320
- #define MMQ_Y_Q5_1 32
319
+ #define MOE_X_Q5_1 4
320
+ #define MOE_Y_Q5_1 32
321
#define NWARPS_Q5_1 4
322
#endif
323
@@ -332,8 +332,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q5_1, 2)
332
const int exp_stride, const int ncols_x, const int nrows_x,
333
const int ncols_y, const int nrows_y, const int nrows_dst,
334
const int top_k) {
335
- const int mmq_x = MMQ_X_Q5_1;
336
- const int mmq_y = MMQ_Y_Q5_1;
335
+ const int mmq_x = MOE_X_Q5_1;
336
+ const int mmq_y = MOE_Y_Q5_1;
337
const int nwarps = NWARPS_Q5_1;
338
339
moe_q<scalar_t, QK5_1, QR5_1, QI5_1, true, block_q5_1, mmq_x, mmq_y, nwarps,
@@ -350,8 +350,8 @@ static void ggml_moe_q5_1_q8_1_cuda(
350
const int exp_stride, const int ncols_x, const int nrows_x,
351
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
352
const int tokens_post_padded, cudaStream_t stream) {
353
- const int mmq_x = MMQ_X_Q5_1;
354
- const int mmq_y = MMQ_Y_Q5_1;
353
+ const int mmq_x = MOE_X_Q5_1;
354
+ const int mmq_y = MOE_Y_Q5_1;
355
const int nwarps = NWARPS_Q5_1;
356
357
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -373,12 +373,12 @@ static void ggml_moe_q5_1_q8_1_cuda(
373
}
374
375
#if defined(USE_ROCM)
376
- #define MMQ_X_Q8_0 64
377
- #define MMQ_Y_Q8_0 128
376
+ #define MOE_X_Q8_0 64
377
+ #define MOE_Y_Q8_0 128
378
#define NWARPS_Q8_0 8
379
#else
380
- #define MMQ_X_Q8_0 4
381
- #define MMQ_Y_Q8_0 32
380
+ #define MOE_X_Q8_0 4
381
+ #define MOE_Y_Q8_0 32
382
#define NWARPS_Q8_0 4
383
#endif
384
@@ -393,8 +393,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q8_0, 2)
393
const int exp_stride, const int ncols_x, const int nrows_x,
394
const int ncols_y, const int nrows_y, const int nrows_dst,
395
const int top_k) {
396
- const int mmq_x = MMQ_X_Q8_0;
397
- const int mmq_y = MMQ_Y_Q8_0;
396
+ const int mmq_x = MOE_X_Q8_0;
397
+ const int mmq_y = MOE_Y_Q8_0;
398
const int nwarps = NWARPS_Q8_0;
399
400
moe_q<scalar_t, QK8_0, QR8_0, QI8_0, false, block_q8_0, mmq_x, mmq_y, nwarps,
@@ -411,8 +411,8 @@ static void ggml_moe_q8_0_q8_1_cuda(
411
const int exp_stride, const int ncols_x, const int nrows_x,
412
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
413
const int tokens_post_padded, cudaStream_t stream) {
414
- const int mmq_x = MMQ_X_Q8_0;
415
- const int mmq_y = MMQ_Y_Q8_0;
414
+ const int mmq_x = MOE_X_Q8_0;
415
+ const int mmq_y = MOE_Y_Q8_0;
416
const int nwarps = NWARPS_Q8_0;
417
418
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -434,12 +434,12 @@ static void ggml_moe_q8_0_q8_1_cuda(
434
}
435
436
#if defined(USE_ROCM)
437
- #define MMQ_X_Q2_K 64
438
- #define MMQ_Y_Q2_K 128
437
+ #define MOE_X_Q2_K 64
438
+ #define MOE_Y_Q2_K 128
439
#define NWARPS_Q2_K 8
440
#else
441
- #define MMQ_X_Q2_K 4
442
- #define MMQ_Y_Q2_K 32
441
+ #define MOE_X_Q2_K 4
442
+ #define MOE_Y_Q2_K 32
443
#define NWARPS_Q2_K 4
444
#endif
445
@@ -454,8 +454,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q2_K, 2)
454
const int exp_stride, const int ncols_x, const int nrows_x,
455
const int ncols_y, const int nrows_y, const int nrows_dst,
456
const int top_k) {
457
- const int mmq_x = MMQ_X_Q2_K;
458
- const int mmq_y = MMQ_Y_Q2_K;
457
+ const int mmq_x = MOE_X_Q2_K;
458
+ const int mmq_y = MOE_Y_Q2_K;
459
const int nwarps = NWARPS_Q2_K;
460
461
moe_q<scalar_t, QK_K, QR2_K, QI2_K, false, block_q2_K, mmq_x, mmq_y, nwarps,
@@ -472,8 +472,8 @@ static void ggml_moe_q2_K_q8_1_cuda(
472
const int exp_stride, const int ncols_x, const int nrows_x,
473
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
474
const int tokens_post_padded, cudaStream_t stream) {
475
- const int mmq_x = MMQ_X_Q2_K;
476
- const int mmq_y = MMQ_Y_Q2_K;
475
+ const int mmq_x = MOE_X_Q2_K;
476
+ const int mmq_y = MOE_Y_Q2_K;
477
const int nwarps = NWARPS_Q2_K;
478
479
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -495,12 +495,12 @@ static void ggml_moe_q2_K_q8_1_cuda(
495
}
496
497
#if defined(USE_ROCM)
498
- #define MMQ_X_Q3_K 64
499
- #define MMQ_Y_Q3_K 128
498
+ #define MOE_X_Q3_K 64
499
+ #define MOE_Y_Q3_K 128
500
#define NWARPS_Q3_K 8
501
#else
502
- #define MMQ_X_Q3_K 4
503
- #define MMQ_Y_Q3_K 32
502
+ #define MOE_X_Q3_K 4
503
+ #define MOE_Y_Q3_K 32
504
#define NWARPS_Q3_K 4
505
#endif
506
@@ -516,8 +516,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q3_K, 2)
516
const int ncols_y, const int nrows_y, const int nrows_dst,
517
const int top_k) {
518
519
- const int mmq_x = MMQ_X_Q3_K;
520
- const int mmq_y = MMQ_Y_Q3_K;
519
+ const int mmq_x = MOE_X_Q3_K;
520
+ const int mmq_y = MOE_Y_Q3_K;
521
const int nwarps = NWARPS_Q3_K;
522
523
moe_q<scalar_t, QK_K, QR3_K, QI3_K, false, block_q3_K, mmq_x, mmq_y, nwarps,
@@ -533,8 +533,8 @@ static void ggml_moe_q3_K_q8_1_cuda(
533
const int exp_stride, const int ncols_x, const int nrows_x,
534
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
535
const int tokens_post_padded, cudaStream_t stream) {
536
- const int mmq_x = MMQ_X_Q3_K;
537
- const int mmq_y = MMQ_Y_Q3_K;
536
+ const int mmq_x = MOE_X_Q3_K;
537
+ const int mmq_y = MOE_Y_Q3_K;
538
const int nwarps = NWARPS_Q3_K;
539
540
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -556,12 +556,12 @@ static void ggml_moe_q3_K_q8_1_cuda(
556
}
557
558
#if defined(USE_ROCM)
559
- #define MMQ_X_Q4_K 64
560
- #define MMQ_Y_Q4_K 128
559
+ #define MOE_X_Q4_K 64
560
+ #define MOE_Y_Q4_K 128
561
#define NWARPS_Q4_K 8
562
#else
563
- #define MMQ_X_Q4_K 4
564
- #define MMQ_Y_Q4_K 32
563
+ #define MOE_X_Q4_K 4
564
+ #define MOE_Y_Q4_K 32
565
#define NWARPS_Q4_K 4
566
#endif
567
@@ -576,8 +576,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q4_K, 2)
576
const int exp_stride, const int ncols_x, const int nrows_x,
577
const int ncols_y, const int nrows_y, const int nrows_dst,
578
const int top_k) {
579
- const int mmq_x = MMQ_X_Q4_K;
580
- const int mmq_y = MMQ_Y_Q4_K;
579
+ const int mmq_x = MOE_X_Q4_K;
580
+ const int mmq_y = MOE_Y_Q4_K;
581
const int nwarps = NWARPS_Q4_K;
582
583
moe_q<scalar_t, QK_K, QR4_K, QI4_K, true, block_q4_K, mmq_x, mmq_y, nwarps,
@@ -594,8 +594,8 @@ static void ggml_moe_q4_K_q8_1_cuda(
594
const int exp_stride, const int ncols_x, const int nrows_x,
595
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
596
const int tokens_post_padded, cudaStream_t stream) {
597
- const int mmq_x = MMQ_X_Q4_K;
598
- const int mmq_y = MMQ_Y_Q4_K;
597
+ const int mmq_x = MOE_X_Q4_K;
598
+ const int mmq_y = MOE_Y_Q4_K;
599
const int nwarps = NWARPS_Q4_K;
600
601
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -617,12 +617,12 @@ static void ggml_moe_q4_K_q8_1_cuda(
617
}
618
619
#if defined(USE_ROCM)
620
- #define MMQ_X_Q5_K 64
621
- #define MMQ_Y_Q5_K 128
620
+ #define MOE_X_Q5_K 64
621
+ #define MOE_Y_Q5_K 128
622
#define NWARPS_Q5_K 8
623
#else
624
- #define MMQ_X_Q5_K 4
625
- #define MMQ_Y_Q5_K 32
624
+ #define MOE_X_Q5_K 4
625
+ #define MOE_Y_Q5_K 32
626
#define NWARPS_Q5_K 4
627
#endif
628
@@ -637,8 +637,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q5_K, 2)
637
const int exp_stride, const int ncols_x, const int nrows_x,
638
const int ncols_y, const int nrows_y, const int nrows_dst,
639
const int top_k) {
640
- const int mmq_x = MMQ_X_Q5_K;
641
- const int mmq_y = MMQ_Y_Q5_K;
640
+ const int mmq_x = MOE_X_Q5_K;
641
+ const int mmq_y = MOE_Y_Q5_K;
642
const int nwarps = NWARPS_Q5_K;
643
644
moe_q<scalar_t, QK_K, QR5_K, QI5_K, true, block_q5_K, mmq_x, mmq_y, nwarps,
@@ -655,8 +655,8 @@ static void ggml_moe_q5_K_q8_1_cuda(
655
const int exp_stride, const int ncols_x, const int nrows_x,
656
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
657
const int tokens_post_padded, cudaStream_t stream) {
658
- const int mmq_x = MMQ_X_Q5_K;
659
- const int mmq_y = MMQ_Y_Q5_K;
658
+ const int mmq_x = MOE_X_Q5_K;
659
+ const int mmq_y = MOE_Y_Q5_K;
660
const int nwarps = NWARPS_Q5_K;
661
662
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;
@@ -678,12 +678,12 @@ static void ggml_moe_q5_K_q8_1_cuda(
678
}
679
680
#if defined(USE_ROCM)
681
- #define MMQ_X_Q6_K 64
682
- #define MMQ_Y_Q6_K 128
681
+ #define MOE_X_Q6_K 64
682
+ #define MOE_Y_Q6_K 128
683
#define NWARPS_Q6_K 8
684
#else
685
- #define MMQ_X_Q6_K 4
686
- #define MMQ_Y_Q6_K 32
685
+ #define MOE_X_Q6_K 4
686
+ #define MOE_Y_Q6_K 32
687
#define NWARPS_Q6_K 4
688
#endif
689
@@ -698,8 +698,8 @@ __launch_bounds__(WARP_SIZE_GGUF* NWARPS_Q6_K, 2)
698
const int exp_stride, const int ncols_x, const int nrows_x,
699
const int ncols_y, const int nrows_y, const int nrows_dst,
700
const int top_k) {
701
- const int mmq_x = MMQ_X_Q6_K;
702
- const int mmq_y = MMQ_Y_Q6_K;
701
+ const int mmq_x = MOE_X_Q6_K;
702
+ const int mmq_y = MOE_Y_Q6_K;
703
const int nwarps = NWARPS_Q6_K;
704
705
moe_q<scalar_t, QK_K, QR6_K, QI6_K, false, block_q6_K, mmq_x, mmq_y, nwarps,
@@ -716,8 +716,8 @@ static void ggml_moe_q6_K_q8_1_cuda(
716
const int exp_stride, const int ncols_x, const int nrows_x,
717
const int ncols_y, const int nrows_y, const int nrows_dst, const int top_k,
718
const int tokens_post_padded, cudaStream_t stream) {
719
- const int mmq_x = MMQ_X_Q6_K;
720
- const int mmq_y = MMQ_Y_Q6_K;
719
+ const int mmq_x = MOE_X_Q6_K;
720
+ const int mmq_y = MOE_Y_Q6_K;
721
const int nwarps = NWARPS_Q6_K;
722
723
const int block_num_x = (nrows_x + mmq_y - 1) / mmq_y;