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, 19:45 UTC
[BugFix] Illegal Memory Access in the blockwise cutlass fp8 GEMMs (#14396)
Lucas Wilkinson committed
Mar 7, 2025 at 00:56 UTC
e5e03c2c1b3708640b913a0bcb9b6815f047e7fa
1 file changed
+1
-1
csrc/cutlass_extensions/gemm/collective/sm90_mma_tma_gmma_ss_warpspecialized_fp8_blockwise_scaling.hpp
+1
-1
@@ -402,7 +402,7 @@ struct CollectiveMma<
402
403
// TODO: test `scale_copy_a` with `ScaleMsPerTile` < 128
404
TiledCopy scale_copy_a = make_tiled_copy(SmemBlockScalingCopyAtomA{},
405
- Layout<Shape<_32, _1>>{}, Layout<Shape<_4, _1>>{}); // (1,1,1)
405
+ Layout<Shape<_32>>{}, Layout<Shape<_1>>{}); // (1,1,1)
406
TiledCopy scale_copy_b = make_tiled_copy(SmemBlockScalingCopyAtomB{},
407
Layout<Shape<_1>>{}, Layout<Shape<_1>>{}); // (1,1,1)
408
ThrCopy thr_scale_copy_a = scale_copy_a.get_slice(threadIdx.x);