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 15, 02:15 UTC
[ROCm] warpSize is being made non constexpr in ROCm 7.0 (#20330)
Signed-off-by: Gregory Shtrasberg <Gregory.Shtrasberg@amd.com>
Gregory Shtrasberg committed
Jul 15, 2025 at 14:01 UTC
ed10f3cea199a7a1f3532fbe367f5c5479a6cae9
4 files changed
+6
-24
csrc/attention/attention_kernels.cuh
+1
-7
@@ -24,6 +24,7 @@
24
25
#include "attention_dtypes.h"
26
#include "attention_utils.cuh"
27
+#include "cuda_compat.h"
28
29
#ifdef USE_ROCM
30
#include <hip/hip_bf16.h>
@@ -33,12 +34,6 @@ typedef __hip_bfloat16 __nv_bfloat16;
34
#include "../quantization/fp8/nvidia/quant_utils.cuh"
35
#endif
36
36
-#ifndef USE_ROCM
37
- #define WARP_SIZE 32
38
-#else
39
- #define WARP_SIZE warpSize
40
-#endif
41
-
37
#define MAX(a, b) ((a) > (b) ? (a) : (b))
38
#define MIN(a, b) ((a) < (b) ? (a) : (b))
39
#define DIVIDE_ROUND_UP(a, b) (((a) + (b) - 1) / (b))
@@ -670,7 +665,6 @@ __global__ void paged_attention_v2_reduce_kernel(
665
666
} // namespace vllm
667
673
-#undef WARP_SIZE
668
#undef MAX
669
#undef MIN
670
#undef DIVIDE_ROUND_UP
csrc/attention/paged_attention_v1.cu
+1
-7
@@ -18,12 +18,7 @@
18
*/
19
20
#include "attention_kernels.cuh"
21
-
22
-#ifndef USE_ROCM
23
- #define WARP_SIZE 32
24
-#else
25
- #define WARP_SIZE warpSize
26
-#endif
21
+#include "cuda_compat.h"
22
23
#define MAX(a, b) ((a) > (b) ? (a) : (b))
24
#define MIN(a, b) ((a) < (b) ? (a) : (b))
@@ -187,7 +182,6 @@ void paged_attention_v1(
182
CALL_V1_LAUNCHER_BLOCK_SIZE)
183
}
184
190
-#undef WARP_SIZE
185
#undef MAX
186
#undef MIN
187
#undef DIVIDE_ROUND_UP
csrc/attention/paged_attention_v2.cu
+1
-7
@@ -18,12 +18,7 @@
18
*/
19
20
#include "attention_kernels.cuh"
21
-
22
-#ifndef USE_ROCM
23
- #define WARP_SIZE 32
24
-#else
25
- #define WARP_SIZE warpSize
26
-#endif
21
+#include "cuda_compat.h"
22
23
#define MAX(a, b) ((a) > (b) ? (a) : (b))
24
#define MIN(a, b) ((a) < (b) ? (a) : (b))
@@ -197,7 +192,6 @@ void paged_attention_v2(
192
CALL_V2_LAUNCHER_BLOCK_SIZE)
193
}
194
200
-#undef WARP_SIZE
195
#undef MAX
196
#undef MIN
197
#undef DIVIDE_ROUND_UP
csrc/cuda_compat.h
+3
-3
@@ -4,10 +4,10 @@
4
#include <hip/hip_runtime.h>
5
#endif
6
7
-#ifndef USE_ROCM
8
- #define WARP_SIZE 32
7
+#if defined(USE_ROCM) && defined(__GFX9__)
8
+ #define WARP_SIZE 64
9
#else
10
- #define WARP_SIZE warpSize
10
+ #define WARP_SIZE 32
11
#endif
12
13
#ifndef USE_ROCM