chenchenghao/sglang
1
0
Fork 0
Code Issues Pull Requests Actions 12 Packages Projects Releases Wiki Activity

fix(sgl-kernel): use >= 120 for SM12x CUDA kernel dispatch (#18750)

Browse Source 
Co-authored-by: Claude Opus 4.6 <noreply@anthropic.com>
This commit is contained in:
blake-snc
2026-02-15 08:44:47 -08:00
committed by GitHub
parent 5fc328465a
commit 0d30896015
3 changed files with 3 additions and 3 deletions
Download Patch File Download Diff File Expand all files Collapse all files

2
sgl-kernel/csrc/gemm/fp8_blockwise_gemm_kernel.cu
View File

@@ -448,7 +448,7 @@ torch::Tensor fp8_blockwise_scaled_mm(
#if defined(CUTLASS_ARCH_MMA_SM120A_SUPPORTED) || defined(CUTLASS_ARCH_MMA_SM120_SUPPORTED)
#if defined(CUDA_VERSION) && CUDA_VERSION >= 12080
if (sm_version == 120) {
if (sm_version >= 120) {
if (out_dtype == torch::kBFloat16) {
sm120_fp8_blockwise_dispatch_shape<cutlass::bfloat16_t>(
out_padded, mat_a_padded, mat_b, scales_a_padded, scales_b);

2
sgl-kernel/csrc/gemm/nvfp4_scaled_mm_kernels.cu
View File

@@ -663,7 +663,7 @@ void cutlass_scaled_fp4_mm_sm100a_sm120a(
// Check SM version and dispatch accordingly
auto sm_version = getSMVersion();
if (sm_version == 120) {
if (sm_version >= 120) {
// Use SM120 specific dispatch
if (out_dtype == at::ScalarType::Half) {
cutlass_fp4_f16_gemm_dispatch_sm120(D, A, B, A_sf, B_sf, alpha, m, n, k, stream);

2
sgl-kernel/csrc/moe/nvfp4_blockwise_moe.cu
View File

@@ -665,7 +665,7 @@ void cutlass_fp4_group_mm(
N,
K);
}
} else if (sm_version == 120) {
} else if (sm_version >= 120) {
if (output.scalar_type() == torch::kBFloat16) {
run_fp4_blockwise_scaled_group_mm_sm120(
output,