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
Revert "Use Cache Hinting for fused_moe kernel (#15511)" (#15645)
Signed-off-by: Wes Medford <wryanmedford@gmail.com>
Wes committed
Mar 27, 2025 at 20:45 UTC
4ae17bf1e242d18d5cbf2eacdaf60185957d6f5b
1 file changed
+4
-12
vllm/model_executor/layers/fused_moe/fused_moe.py
+4
-12
@@ -189,11 +189,7 @@ def fused_moe_kernel_gptq_awq(
189
mask=token_mask[:, None] &
190
(offs_k[None, :] < K - k * BLOCK_SIZE_K),
191
other=0.0)
192
- b = tl.load(
193
- b_ptrs,
194
- cache_modifier=".cg",
195
- eviction_policy="evict_last",
196
- )
192
+ b = tl.load(b_ptrs)
193
if use_int4_w4a16:
194
b = (b >> b_shifter) & 0xF
195
@@ -395,13 +391,9 @@ def fused_moe_kernel(
391
mask=token_mask[:, None] &
392
(offs_k[None, :] < K - k * BLOCK_SIZE_K),
393
other=0.0)
398
- b = tl.load(
399
- b_ptrs,
400
- mask=offs_k[:, None] < K - k * BLOCK_SIZE_K,
401
- other=0.0,
402
- cache_modifier=".cg",
403
- eviction_policy="evict_last",
404
- )
394
+ b = tl.load(b_ptrs,
395
+ mask=offs_k[:, None] < K - k * BLOCK_SIZE_K,
396
+ other=0.0)
397
# We accumulate along the K dimension.
398
if use_int8_w8a16:
399
accumulator = tl.dot(a, b.to(compute_type), acc=accumulator)