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, 09:45 UTC
fix CUDA_check redefinition in #17918 (#18287)
Signed-off-by: Lucia Fang <fanglu@fb.com> Co-authored-by: Lucia (Lu) Fang <fanglu@meta.com>
Lucia Fang committed
May 19, 2025 at 13:42 UTC
258bf621d5e533b01026b73fbfb31b746f68684f
2 files changed
+5
-12
csrc/cutlass_extensions/common.hpp
-9
@@ -15,15 +15,6 @@
15
cutlassGetStatusString(error)); \
16
}
17
18
-/**
19
- * Panic wrapper for unwinding CUDA runtime errors
20
- */
21
-#define CUDA_CHECK(status) \
22
- { \
23
- cudaError_t error = status; \
24
- TORCH_CHECK(error == cudaSuccess, cudaGetErrorString(error)); \
25
- }
26
-
18
inline int get_cuda_max_shared_memory_per_block_opt_in(int const device) {
19
int max_shared_mem_per_block_opt_in = 0;
20
cudaDeviceGetAttribute(&max_shared_mem_per_block_opt_in,
csrc/sparse/cutlass/sparse_scaled_mm_c3x.cuh
+5
-3
@@ -8,6 +8,8 @@
8
9
#include <ATen/cuda/CUDAContext.h>
10
11
+#include "cuda_utils.h"
12
+
13
#include "cutlass/cutlass.h"
14
15
#include "cutlass/gemm/device/gemm_universal_adapter.h"
@@ -95,9 +97,9 @@ struct cutlass_sparse_3x_gemm {
97
// clang-format off
98
using CollectiveMainloop =
99
typename cutlass::gemm::collective::CollectiveBuilder<
98
- cutlass::arch::Sm90, cutlass::arch::OpClassSparseTensorOp,
99
- ElementAB, cutlass::layout::RowMajor, AlignmentAB,
100
- ElementAB, cutlass::layout::ColumnMajor, AlignmentAB,
100
+ cutlass::arch::Sm90, cutlass::arch::OpClassSparseTensorOp,
101
+ ElementAB, cutlass::layout::RowMajor, AlignmentAB,
102
+ ElementAB, cutlass::layout::ColumnMajor, AlignmentAB,
103
ElementAcc, TileShape, ClusterShape,
104
Stages,
105
KernelSchedule>::CollectiveOp;