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 4, 12:30 UTC
[CI/Build] Enable FP8 on NVIDIA Thor (#39712)
Signed-off-by: DarkLight1337 <tlleungac@connect.ust.hk>
Cyrus Leung committed
Apr 30, 2026 at 00:48 UTC
4a42aba380bf9cac47009a7307a9d91dd2222d84
4 files changed
+13
-38
csrc/cutlass_extensions/common.hpp
+10
-35
@@ -96,44 +96,14 @@ struct enable_sm90_or_later : Kernel {
96
};
97
98
template <typename Kernel>
99
-struct enable_sm90_only : Kernel {
99
+struct enable_sm100_to_sm120 : Kernel {
100
template <typename... Args>
101
CUTLASS_DEVICE void operator()(Args&&... args) {
102
#if defined __CUDA_ARCH__
103
- #if __CUDA_ARCH__ == 900
103
+ #if (__CUDA_ARCH__ >= 1000 && __CUDA_ARCH__ < 1200)
104
Kernel::operator()(std::forward<Args>(args)...);
105
#else
106
- printf("This kernel only supports sm90.\n");
107
- asm("trap;");
108
- #endif
109
-#endif
110
- }
111
-};
112
-
113
-template <typename Kernel>
114
-struct enable_sm100f_only : Kernel {
115
- template <typename... Args>
116
- CUTLASS_DEVICE void operator()(Args&&... args) {
117
-#if defined __CUDA_ARCH__
118
- #if __CUDA_ARCH__ == 1000 || __CUDA_ARCH__ == 1030
119
- Kernel::operator()(std::forward<Args>(args)...);
120
- #else
121
- printf("This kernel only supports sm100f.\n");
122
- asm("trap;");
123
- #endif
124
-#endif
125
- }
126
-};
127
-
128
-template <typename Kernel>
129
-struct enable_sm100a_only : Kernel {
130
- template <typename... Args>
131
- CUTLASS_DEVICE void operator()(Args&&... args) {
132
-#if defined __CUDA_ARCH__
133
- #if __CUDA_ARCH__ == 1000
134
- Kernel::operator()(std::forward<Args>(args)...);
135
- #else
136
- printf("This kernel only supports sm100a.\n");
106
+ printf("This kernel only supports sm[100, 120).\n");
107
asm("trap;");
108
#endif
109
#endif
@@ -148,7 +118,7 @@ struct enable_sm120_only : Kernel {
118
#if __CUDA_ARCH__ == 1200
119
Kernel::operator()(std::forward<Args>(args)...);
120
#else
151
- printf("This kernel only supports sm120.\n");
121
+ printf("This kernel only supports sm120a.\n");
122
asm("trap;");
123
#endif
124
#endif
@@ -160,8 +130,13 @@ template <typename Kernel>
130
struct enable_sm120_family : Kernel {
131
template <typename... Args>
132
CUTLASS_DEVICE void operator()(Args&&... args) {
163
-#if defined __CUDA_ARCH__ && (__CUDA_ARCH__ >= 1200 && __CUDA_ARCH__ < 1300)
133
+#if defined __CUDA_ARCH__
134
+ #if (__CUDA_ARCH__ >= 1200 && __CUDA_ARCH__ < 1300)
135
Kernel::operator()(std::forward<Args>(args)...);
136
+ #else
137
+ printf("This kernel only supports sm120f.\n");
138
+ asm("trap;");
139
+ #endif
140
#endif
141
}
142
};
csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm.cuh
+1
-1
@@ -141,7 +141,7 @@ struct cutlass_3x_gemm_sm100 {
141
sizeof(typename CollectiveEpilogue::SharedStorage))>,
142
KernelSchedule>::CollectiveOp;
143
144
- using GemmKernel = enable_sm100f_only<cutlass::gemm::kernel::GemmUniversal<
144
+ using GemmKernel = enable_sm100_to_sm120<cutlass::gemm::kernel::GemmUniversal<
145
Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>>;
146
};
147
csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_blockwise_sm100_fp8_dispatch.cuh
+1
-1
@@ -125,7 +125,7 @@ struct cutlass_3x_gemm_fp8_blockwise {
125
MainloopScheduler
126
>::CollectiveOp>;
127
128
- using KernelType = enable_sm100f_only<cutlass::gemm::kernel::GemmUniversal<
128
+ using KernelType = enable_sm100_to_sm120<cutlass::gemm::kernel::GemmUniversal<
129
Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue>>;
130
131
struct GemmKernel : public KernelType {};
csrc/libtorch_stable/quantization/w8a8/cutlass/c3x/scaled_mm_sm100_fp8_dispatch.cuh
+1
-1
@@ -92,7 +92,7 @@ struct cutlass_3x_gemm_sm100_fp8 {
92
// -----------------------------------------------------------
93
// Kernel definition
94
// -----------------------------------------------------------
95
- using GemmKernel = enable_sm100f_only<cutlass::gemm::kernel::GemmUniversal<
95
+ using GemmKernel = enable_sm100_to_sm120<cutlass::gemm::kernel::GemmUniversal<
96
Shape<int, int, int, int>, CollectiveMainloop, CollectiveEpilogue, void>>;
97
};
98