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
[Compile] Fix Compile Warning for Ignoring `MIN_BLOCK_PER_SM` (#25193)
Signed-off-by: yewentao256 <zhyanwentao@126.com>
Wentao Ye committed
Sep 19, 2025 at 18:23 UTC
711e912946d23f4ccc1f554b1524c960553c5e28
4 files changed
+53
-9
csrc/launch_bounds_utils.h
new
+38
@@ -0,0 +1,38 @@
1
+#pragma once
2
+
3
+#include <cuda_runtime_api.h>
4
+#include <algorithm>
5
+
6
+// maximum blocks per SM cap
7
+#ifndef VLLM_LAUNCH_BLOCKS_CAP
8
+ #define VLLM_LAUNCH_BLOCKS_CAP 4
9
+#endif
10
+
11
+// compile-time estimate of max threads per SM for launch bounds.
12
+#ifndef VLLM_MAX_THREADS_PER_SM
13
+ #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 300
14
+ #define VLLM_MAX_THREADS_PER_SM 1536
15
+ #else
16
+ #define VLLM_MAX_THREADS_PER_SM 2048
17
+ #endif
18
+#endif
19
+
20
+// compute the number of blocks per SM to request in __launch_bounds__
21
+#define VLLM_BLOCKS_DIV(VAL) (VLLM_MAX_THREADS_PER_SM / (VAL))
22
+#define VLLM_CLAMP_BLOCKS_PER_SM(VAL) \
23
+ (((VAL) <= 0) \
24
+ ? 1 \
25
+ : (((VAL) < VLLM_LAUNCH_BLOCKS_CAP) ? (VAL) : VLLM_LAUNCH_BLOCKS_CAP))
26
+#define VLLM_BLOCKS_PER_SM(BLOCK_THREADS) \
27
+ VLLM_CLAMP_BLOCKS_PER_SM(VLLM_BLOCKS_DIV(BLOCK_THREADS))
28
+
29
+// runtime-time helper to compute blocks/SM
30
+static inline int vllm_runtime_blocks_per_sm(int block_threads) {
31
+ int device = -1;
32
+ cudaGetDevice(&device);
33
+ int max_threads_per_sm = VLLM_MAX_THREADS_PER_SM;
34
+ cudaDeviceGetAttribute(&max_threads_per_sm,
35
+ cudaDevAttrMaxThreadsPerMultiProcessor, device);
36
+ int blocks = (block_threads > 0) ? (max_threads_per_sm / block_threads) : 1;
37
+ return VLLM_CLAMP_BLOCKS_PER_SM(blocks);
38
+}
csrc/quantization/fp4/activation_nvfp4_quant_fusion_kernels.cu
+4
-2
@@ -26,6 +26,7 @@
26
#include "dispatch_utils.h"
27
28
#include "cuda_utils.h"
29
+#include "launch_bounds_utils.h"
30
#include "nvfp4_utils.cuh"
31
32
namespace vllm {
@@ -63,7 +64,7 @@ __inline__ __device__ PackedVec<Type> compute_silu_mul(PackedVec<Type>& vec,
64
65
// Use UE4M3 by default.
66
template <class Type, bool UE8M0_SF = false>
66
-__global__ void __launch_bounds__(1024, 4)
67
+__global__ void __launch_bounds__(1024, VLLM_BLOCKS_PER_SM(1024))
68
silu_mul_cvt_fp16_to_fp4(int32_t numRows, int32_t numCols, Type const* in,
69
float const* SFScale, uint32_t* out,
70
uint32_t* SFout) {
@@ -131,7 +132,8 @@ void silu_and_mul_nvfp4_quant_sm1xxa(torch::Tensor& output, // [..., d]
132
const at::cuda::OptionalCUDAGuard device_guard(device_of(input));
133
auto stream = at::cuda::getCurrentCUDAStream(input.get_device());
134
dim3 block(std::min(int(n / ELTS_PER_THREAD), 1024));
134
- int const numBlocksPerSM = 2048 / block.x;
135
+ int const numBlocksPerSM =
136
+ vllm_runtime_blocks_per_sm(static_cast<int>(block.x));
137
dim3 grid(std::min(int(m), multiProcessorCount * numBlocksPerSM));
138
139
VLLM_DISPATCH_HALF_TYPES(
csrc/quantization/fp4/nvfp4_experts_quant.cu
+6
-4
@@ -26,12 +26,13 @@
26
#include "dispatch_utils.h"
27
28
#include "nvfp4_utils.cuh"
29
+#include "launch_bounds_utils.h"
30
31
namespace vllm {
32
33
// Use UE4M3 by default.
34
template <class Type, bool UE8M0_SF = false, bool SMALL_NUM_EXPERTS = false>
34
-__global__ void __launch_bounds__(512, 4)
35
+__global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
36
cvt_fp16_to_fp4(int32_t numRows, int32_t numCols, Type const* in,
37
float const* SFScale, uint32_t* out, uint32_t* SFout,
38
uint32_t* input_offset_by_experts,
@@ -129,7 +130,7 @@ __global__ void __launch_bounds__(512, 4)
130
131
// Kernel for LARGE_M_TOPK = true (large m_topk optimized version)
132
template <class Type, bool UE8M0_SF = false, bool SMALL_NUM_EXPERTS = false>
132
-__global__ void __launch_bounds__(1024, 4)
133
+__global__ void __launch_bounds__(1024, VLLM_BLOCKS_PER_SM(1024))
134
cvt_fp16_to_fp4(int32_t numRows, int32_t numCols, Type const* in,
135
float const* SFScale, uint32_t* out, uint32_t* SFout,
136
uint32_t* input_offset_by_experts,
@@ -233,8 +234,9 @@ void quant_impl(void* output, void* output_scale, void* input,
234
int const workSizePerRow = k / ELTS_PER_THREAD;
235
int const totalWorkSize = m_topk * workSizePerRow;
236
dim3 block(std::min(workSizePerRow, 512));
236
- // Get number of blocks per SM (assume we can fully utilize the SM).
237
- int const numBlocksPerSM = 2048 / block.x;
237
+ // Get number of blocks per SM
238
+ int const numBlocksPerSM =
239
+ vllm_runtime_blocks_per_sm(static_cast<int>(block.x));
240
dim3 grid(std::min(static_cast<int>((totalWorkSize + block.x - 1) / block.x),
241
multiProcessorCount * numBlocksPerSM));
242
while (grid.x <= multiProcessorCount && block.x > 64) {
csrc/quantization/fp4/nvfp4_quant_kernels.cu
+5
-3
@@ -26,13 +26,14 @@
26
#include "dispatch_utils.h"
27
28
#include "cuda_utils.h"
29
+#include "launch_bounds_utils.h"
30
#include "nvfp4_utils.cuh"
31
32
namespace vllm {
33
34
// Use UE4M3 by default.
35
template <class Type, bool UE8M0_SF = false>
35
-__global__ void __launch_bounds__(512, 4)
36
+__global__ void __launch_bounds__(512, VLLM_BLOCKS_PER_SM(512))
37
cvt_fp16_to_fp4(int32_t numRows, int32_t numCols, Type const* in,
38
float const* SFScale, uint32_t* out, uint32_t* SFout) {
39
using PackedVec = PackedVec<Type>;
@@ -75,8 +76,9 @@ void invokeFP4Quantization(int m, int n, T const* input, float const* SFScale,
76
// Grid, Block size.
77
// Each thread converts 8 values.
78
dim3 block(std::min(int(n / ELTS_PER_THREAD), 512));
78
- // Get number of blocks per SM (assume we can fully utilize the SM).
79
- int const numBlocksPerSM = 2048 / block.x;
79
+ // Get number of blocks per SM
80
+ int const numBlocksPerSM =
81
+ vllm_runtime_blocks_per_sm(static_cast<int>(block.x));
82
dim3 grid(std::min(int(m), multiProcessorCount * numBlocksPerSM));
83
84
// Launch the cvt kernel.