From e95668abc7811c53e427ae7f986d0e9bc8af6bfd Mon Sep 17 00:00:00 2001 From: Serge Panev Date: Thu, 22 Jan 2026 05:21:11 +0100 Subject: [PATCH] [NVIDIA] Fix CUDA arch requirement in nvfp4 cast (#12581) Signed-off-by: Serge Panev Co-authored-by: Fan Yin <1106310035@qq.com> --- sgl-kernel/csrc/gemm/nvfp4_quant.cuh | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sgl-kernel/csrc/gemm/nvfp4_quant.cuh b/sgl-kernel/csrc/gemm/nvfp4_quant.cuh index bc7c612ad..82f22f479 100644 --- a/sgl-kernel/csrc/gemm/nvfp4_quant.cuh +++ b/sgl-kernel/csrc/gemm/nvfp4_quant.cuh @@ -52,7 +52,7 @@ constexpr int CVT_FP4_SF_VEC_SIZE = 16; inline __device__ uint32_t fp32_vec_to_e2m1(float (&array)[8]) { // PTX instructions used here requires >= sm100f. #if CUTLASS_ARCH_MMA_SM100A_ENABLED || CUTLASS_ARCH_MMA_SM103A_ENABLED || CUTLASS_ARCH_MMA_SM120A_ENABLED || \ - (defined(__CUDA_ARCH_FAMILY_SPECIFIC__) && (__CUDA_ARCH_FAMILY_SPECIFIC__ > 1000)) + (defined(__CUDA_ARCH_FAMILY_SPECIFIC__) && (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1000)) uint32_t val; asm volatile( "{\n" @@ -87,7 +87,7 @@ inline __device__ uint32_t fp32_vec_to_e2m1(float (&array)[8]) { inline __device__ uint32_t fp32_vec_to_e2m1(float2 (&array)[4]) { // PTX instructions used here requires >= sm100f. #if CUTLASS_ARCH_MMA_SM100A_ENABLED || CUTLASS_ARCH_MMA_SM103A_ENABLED || CUTLASS_ARCH_MMA_SM120A_ENABLED || \ - (defined(__CUDA_ARCH_FAMILY_SPECIFIC__) && (__CUDA_ARCH_FAMILY_SPECIFIC__ > 1000)) + (defined(__CUDA_ARCH_FAMILY_SPECIFIC__) && (__CUDA_ARCH_FAMILY_SPECIFIC__ >= 1000)) uint32_t val; asm volatile( "{\n"