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 14, 05:30 UTC
silu-v1: Fix EPS not being used during max-reduction (#25069)
Signed-off-by: elvircrn <elvircrn@gmail.com>
Elvir Crnčević committed
Sep 18, 2025 at 12:25 UTC
abdfcd4f3dc21dc162baf6887f658fb0f2f3d783
1 file changed
+1
-2
csrc/quantization/activation_kernels.cu
+1
-2
@@ -365,7 +365,6 @@ __global__ void silu_mul_fp8_quant_deep_gemm_kernel(
365
int32_t compute_pipeline_offset_64 = 0;
366
367
for (int32_t t = n_tokens_lower; t < n_tokens_upper; ++t) {
368
- __nv_bfloat16 y_max_bf16 = EPS;
368
__nv_bfloat162 results_bf162[2];
369
370
cp_async_wait<NUM_STAGES - 2>();
@@ -405,7 +404,7 @@ __global__ void silu_mul_fp8_quant_deep_gemm_kernel(
404
auto _y_max2 =
405
__hmax2(__habs2(results_bf162[0]), __habs2(results_bf162[1]));
406
408
- y_max_bf16 = __hmax(_y_max2.x, _y_max2.y);
407
+ __nv_bfloat16 y_max_bf16 = __hmax(EPS, __hmax(_y_max2.x, _y_max2.y));
408
409
// An entire group is assigned to a single warp, so a simple warp reduce
410
// is used.