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, 14:45 UTC
[BugFix][ROCm] Fix GGUF MoE Dispatch Block_Dim for ROCm (#16247)
Signed-off-by: Tianyuan Wu <Tianyuan.Wu@amd.com>
TY-AMD committed
Apr 8, 2025 at 20:10 UTC
9351f91be96c58167631e43feb807a78cf2f0340
1 file changed
+10
-10
csrc/quantization/gguf/moe.cuh
+10
-10
@@ -129,7 +129,7 @@ static __device__ __forceinline__ void moe_q(
129
}
130
131
#if defined(USE_ROCM)
132
- #define MOE_X_Q4_0 64
132
+ #define MOE_X_Q4_0 8
133
#define MOE_Y_Q4_0 128
134
#define NWARPS_Q4_0 8
135
#else
@@ -190,7 +190,7 @@ static void ggml_moe_q4_0_q8_1_cuda(
190
}
191
192
#if defined(USE_ROCM)
193
- #define MOE_X_Q4_1 64
193
+ #define MOE_X_Q4_1 8
194
#define MOE_Y_Q4_1 128
195
#define NWARPS_Q4_1 8
196
#else
@@ -251,7 +251,7 @@ static void ggml_moe_q4_1_q8_1_cuda(
251
}
252
253
#if defined(USE_ROCM)
254
- #define MOE_X_Q5_0 64
254
+ #define MOE_X_Q5_0 8
255
#define MOE_Y_Q5_0 128
256
#define NWARPS_Q5_0 8
257
#else
@@ -312,7 +312,7 @@ static void ggml_moe_q5_0_q8_1_cuda(
312
}
313
314
#if defined(USE_ROCM)
315
- #define MOE_X_Q5_1 64
315
+ #define MOE_X_Q5_1 8
316
#define MOE_Y_Q5_1 128
317
#define NWARPS_Q5_1 8
318
#else
@@ -373,7 +373,7 @@ static void ggml_moe_q5_1_q8_1_cuda(
373
}
374
375
#if defined(USE_ROCM)
376
- #define MOE_X_Q8_0 64
376
+ #define MOE_X_Q8_0 8
377
#define MOE_Y_Q8_0 128
378
#define NWARPS_Q8_0 8
379
#else
@@ -434,7 +434,7 @@ static void ggml_moe_q8_0_q8_1_cuda(
434
}
435
436
#if defined(USE_ROCM)
437
- #define MOE_X_Q2_K 64
437
+ #define MOE_X_Q2_K 8
438
#define MOE_Y_Q2_K 128
439
#define NWARPS_Q2_K 8
440
#else
@@ -495,7 +495,7 @@ static void ggml_moe_q2_K_q8_1_cuda(
495
}
496
497
#if defined(USE_ROCM)
498
- #define MOE_X_Q3_K 64
498
+ #define MOE_X_Q3_K 8
499
#define MOE_Y_Q3_K 128
500
#define NWARPS_Q3_K 8
501
#else
@@ -556,7 +556,7 @@ static void ggml_moe_q3_K_q8_1_cuda(
556
}
557
558
#if defined(USE_ROCM)
559
- #define MOE_X_Q4_K 64
559
+ #define MOE_X_Q4_K 8
560
#define MOE_Y_Q4_K 128
561
#define NWARPS_Q4_K 8
562
#else
@@ -617,7 +617,7 @@ static void ggml_moe_q4_K_q8_1_cuda(
617
}
618
619
#if defined(USE_ROCM)
620
- #define MOE_X_Q5_K 64
620
+ #define MOE_X_Q5_K 8
621
#define MOE_Y_Q5_K 128
622
#define NWARPS_Q5_K 8
623
#else
@@ -678,7 +678,7 @@ static void ggml_moe_q5_K_q8_1_cuda(
678
}
679
680
#if defined(USE_ROCM)
681
- #define MOE_X_Q6_K 64
681
+ #define MOE_X_Q6_K 8
682
#define MOE_Y_Q6_K 128
683
#define NWARPS_Q6_K 8
684
#else