[Bugfix] Fix accuracy issue for silu_mul + nvfp4 quant fusion kernel

[elvischenv]

Purpose

• Fix accuracy issue for silu_mul + nvfp4 quant fusion kernel
  • In previous [NVIDIA] Support SiluMul + NVFP4 quant fusion #23671, the original author used an approximate way to compute silu. The sigmoid was computed by tanh(x/2) * 0.5 + 0.5, while the standard definition should be 1 / (1 + e^(-x)).
  • We didn't observe the accuracy issues previously by offline lm_eval evaluation. see [NVIDIA] Support SiluMul + NVFP4 quant fusion #23671. However, when doing online lm_eval, the accuracy is dropped slightly.
  • We fixed the issue by using a standard implementation to compute silu_mul.

vllm/csrc/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu

Lines 34 to 54 in 5679399

	__inline__ __device__ PackedVec<Type> compute_silu(PackedVec<Type>& vec,
	PackedVec<Type>& vec2) {
	PackedVec<Type> result;
	#pragma unroll
	for (int i = 0; i < CVT_FP4_ELTS_PER_THREAD / 2; ++i) {
	if constexpr (std::is_same_v<Type, half>) {
	half2 val(0.5f, 0.5f);
	half2 t0 = __hmul2(vec.elts[i], val);
	half2 t1 = __hfma2(h2tanh(t0), val, val);
	half2 t2 = __hmul2(vec.elts[i], t1);
	result.elts[i] = __hmul2(t2, vec2.elts[i]);
	} else {
	__nv_bfloat162 val(0.5f, 0.5f);
	__nv_bfloat162 t0 = __hmul2(vec.elts[i], val);
	__nv_bfloat162 t1 = __hfma2(h2tanh(t0), val, val);
	__nv_bfloat162 t2 = __hmul2(vec.elts[i], t1);
	result.elts[i] = __hmul2(t2, vec2.elts[i]);
	}
	}
	return result;
	}

• Do some cleanups for activation_nvfp4_quant_fusion_kernels.cu. We don't need an extra silu_and_cvt_warp_fp16_to_fp4 but just compute_silu_mul + cvt_warp_fp16_to_fp4(reuse from the nvfp4_utils).
• There were only 3 tests covered the silu_mul + quant fusion in test_silu_mul_quant_fusion.py. Improved the test coverage.
• Clean up the kernel test test_silu_mul_nvfp4_quant.py. Removed lots of unnecessary components.

Test Plan && Test Result

Unit test:

tests/compile/test_silu_mul_quant_fusion.py

====== 24 passed, 8 skipped, 5 warnings in 11.79s =====

tests/kernels/quantization/test_silu_mul_nvfp4_quant.py

======== 8 passed in 1.44s ======

E2E online lm_eval:

main not fused:

local-completions (base_url=http://0.0.0.0:8000/v1/completions,model=nvidia/Llama-3.3-70B-Instruct-FP4,tokenized_requests=False,tokenizer_backend=None,num_concurrent=128,timeout=120,max_retries=5), gen_kwargs: (None), limit: None, num_fewshot: None, batch_size: 1
|Tasks|Version|     Filter     |n-shot|  Metric   |   |Value |   |Stderr|
|-----|------:|----------------|-----:|-----------|---|-----:|---|-----:|
|gsm8k|      3|flexible-extract|     5|exact_match|↑  |0.9257|±  |0.0072|
|     |       |strict-match    |     5|exact_match|↑  |0.6209|±  |0.0134|

main fused(the accuracy dropped slightly):

local-completions (base_url=http://0.0.0.0:8000/v1/completions,model=nvidia/Llama-3.3-70B-Instruct-FP4,tokenized_requests=False,tokenizer_backend=None,num_concurrent=128,timeout=120,max_retries=5), gen_kwargs: (None), limit: None, num_fewshot: None, batch_size: 1
|Tasks|Version|     Filter     |n-shot|  Metric   |   |Value |   |Stderr|
|-----|------:|----------------|-----:|-----------|---|-----:|---|-----:|
|gsm8k|      3|flexible-extract|     5|exact_match|↑  |0.9287|±  |0.0071|
|     |       |strict-match    |     5|exact_match|↑  |0.6035|±  |0.0135|

PR fused:

local-completions (base_url=http://0.0.0.0:8000/v1/completions,model=nvidia/Llama-3.3-70B-Instruct-FP4,tokenized_requests=False,tokenizer_backend=None,num_concurrent=128,timeout=120,max_retries=5), gen_kwargs: (None), limit: None, num_fewshot: None, batch_size: 1
|Tasks|Version|     Filter     |n-shot|  Metric   |   |Value |   |Stderr|
|-----|------:|----------------|-----:|-----------|---|-----:|---|-----:|
|gsm8k|      3|flexible-extract|     5|exact_match|↑  |0.9325|±  |0.0069|
|     |       |strict-match    |     5|exact_match|↑  |0.6255|±  |0.0133|

Perf:

main not fused:

triton_poi_fused_mul_silu_4: 4.288 μs
cvt_fp16_to_fp4: 4.544 μs

main fused:

silu_mul_cvt_fp16_to_fp4: 6.527 μs

PR fused:

silu_mul_cvt_fp16_to_fp4: 6.624 μs

---

Essential Elements of an Effective PR Description Checklist

• The purpose of the PR, such as "Fix some issue (link existing issues this PR will resolve)".
• The test plan, such as providing test command.
• The test results, such as pasting the results comparison before and after, or e2e results
• (Optional) The necessary documentation update, such as updating supported_models.md and examples for a new model.
• (Optional) Release notes update. If your change is user facing, please update the release notes draft in the Google Doc.

[pytorch-bot]

No ciflow labels are configured for this repo.
For information on how to enable CIFlow bot see this wiki

[ProExpertProg]

Can you compare to the Inductor-generated fused kernel?

[mgoin]

Looks good to me as a bugfix, we should merge. @ProExpertProg I don't think we have nvfp4 quant in torch implemented in a serious way, so we should leave that to future work

[yewentao256]

LGTM, thanks for the work!