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, 09:45 UTC
[Bugfix] fix `an illegal memory access was encountered` of marlin kernel + act_order (#18245)
Jinzhen Lin committed
May 17, 2025 at 07:02 UTC
e73b7dfd69ae353711362b58180aa32b86a32095
3 files changed
+29
-22
csrc/moe/marlin_moe_wna16/marlin_template.h
+14
-11
@@ -1767,17 +1767,20 @@ __global__ void Marlin(
1767
1768
if constexpr (has_act_order) {
1769
slice_k_start += tb_k * stages;
1770
- slice_k_start_shared_fetch += tb_k * stages;
1771
- int first_group_id = g_idx[slice_k_start];
1772
- int last_g_idx = slice_k_start + stages * tb_k * 2;
1773
- if (last_g_idx >= prob_k) {
1774
- last_g_idx = prob_k - 1;
1775
- }
1776
- int last_group_id = g_idx[last_g_idx];
1777
- if (last_group_id >= sh_first_group_id + sh_num_groups) {
1778
- fetch_act_order_scales_to_shared(false, first_group_id,
1779
- last_group_id);
1780
- __syncthreads();
1770
+
1771
+ if (slice_k_start < prob_k) {
1772
+ slice_k_start_shared_fetch += tb_k * stages;
1773
+ int first_group_id = g_idx[slice_k_start];
1774
+ int last_g_idx = slice_k_start + stages * tb_k * 2;
1775
+ if (last_g_idx >= prob_k) {
1776
+ last_g_idx = prob_k - 1;
1777
+ }
1778
+ int last_group_id = g_idx[last_g_idx];
1779
+ if (last_group_id >= sh_first_group_id + sh_num_groups) {
1780
+ fetch_act_order_scales_to_shared(false, first_group_id,
1781
+ last_group_id);
1782
+ __syncthreads();
1783
+ }
1784
}
1785
}
1786
if (slice_iters == 0) {
csrc/quantization/gptq_marlin/marlin_template.h
+14
-10
@@ -1588,16 +1588,20 @@ __global__ void Marlin(
1588
1589
if constexpr (has_act_order) {
1590
slice_k_start += tb_k * stages;
1591
- slice_k_start_shared_fetch += tb_k * stages;
1592
- int first_group_id = g_idx[slice_k_start];
1593
- int last_g_idx = slice_k_start + stages * tb_k * 2;
1594
- if (last_g_idx >= prob_k) {
1595
- last_g_idx = prob_k - 1;
1596
- }
1597
- int last_group_id = g_idx[last_g_idx];
1598
- if (last_group_id >= sh_first_group_id + sh_num_groups) {
1599
- fetch_act_order_scales_to_shared(false, first_group_id, last_group_id);
1600
- __syncthreads();
1591
+
1592
+ if (slice_k_start < prob_k) {
1593
+ slice_k_start_shared_fetch += tb_k * stages;
1594
+ int first_group_id = g_idx[slice_k_start];
1595
+ int last_g_idx = slice_k_start + stages * tb_k * 2;
1596
+ if (last_g_idx >= prob_k) {
1597
+ last_g_idx = prob_k - 1;
1598
+ }
1599
+ int last_group_id = g_idx[last_g_idx];
1600
+ if (last_group_id >= sh_first_group_id + sh_num_groups) {
1601
+ fetch_act_order_scales_to_shared(false, first_group_id,
1602
+ last_group_id);
1603
+ __syncthreads();
1604
+ }
1605
}
1606
}
1607
tests/weight_loading/models.txt
+1
-1
@@ -2,7 +2,7 @@ gptq_marlin, robertgshaw2/zephyr-7b-beta-channelwise-gptq, main
2
gptq_marlin, TheBloke/Llama-2-7B-GPTQ, main
3
gptq_marlin, TheBloke/TinyLlama-1.1B-Chat-v1.0-GPTQ, main
4
gptq_marlin, TheBloke/TinyLlama-1.1B-Chat-v1.0-GPTQ, gptq-8bit--1g-actorder_True
5
-#gptq_marlin, TheBloke/TinyLlama-1.1B-Chat-v1.0-GPTQ, gptq-8bit-32g-actorder_True
5
+gptq_marlin, TheBloke/TinyLlama-1.1B-Chat-v1.0-GPTQ, gptq-8bit-32g-actorder_True
6
gptq_marlin, TechxGenus/gemma-1.1-2b-it-GPTQ, main
7
gptq, robertgshaw2/zephyr-7b-beta-channelwise-gptq, main
8
gptq, TheBloke/Llama-2-7B-GPTQ, main