From 258bf621d5e533b01026b73fbfb31b746f68684f Mon Sep 17 00:00:00 2001 From: Lucia Fang <116399278+luccafong@users.noreply.github.com> Date: Mon, 19 May 2025 13:42:35 -0700 Subject: [PATCH] fix CUDA_check redefinition in #17918 (#18287) Signed-off-by: Lucia Fang Co-authored-by: Lucia (Lu) Fang --- csrc/cutlass_extensions/common.hpp | 9 --------- csrc/sparse/cutlass/sparse_scaled_mm_c3x.cuh | 8 +++++--- 2 files changed, 5 insertions(+), 12 deletions(-) diff --git a/csrc/cutlass_extensions/common.hpp b/csrc/cutlass_extensions/common.hpp index 0877da52435eb..195872e8edd3e 100644 --- a/csrc/cutlass_extensions/common.hpp +++ b/csrc/cutlass_extensions/common.hpp @@ -15,15 +15,6 @@ cutlassGetStatusString(error)); \ } -/** - * Panic wrapper for unwinding CUDA runtime errors - */ -#define CUDA_CHECK(status) \ - { \ - cudaError_t error = status; \ - TORCH_CHECK(error == cudaSuccess, cudaGetErrorString(error)); \ - } - inline int get_cuda_max_shared_memory_per_block_opt_in(int const device) { int max_shared_mem_per_block_opt_in = 0; cudaDeviceGetAttribute(&max_shared_mem_per_block_opt_in, diff --git a/csrc/sparse/cutlass/sparse_scaled_mm_c3x.cuh b/csrc/sparse/cutlass/sparse_scaled_mm_c3x.cuh index 9c8a50332ad00..c22523da4e439 100644 --- a/csrc/sparse/cutlass/sparse_scaled_mm_c3x.cuh +++ b/csrc/sparse/cutlass/sparse_scaled_mm_c3x.cuh @@ -8,6 +8,8 @@ #include +#include "cuda_utils.h" + #include "cutlass/cutlass.h" #include "cutlass/gemm/device/gemm_universal_adapter.h" @@ -95,9 +97,9 @@ struct cutlass_sparse_3x_gemm { // clang-format off using CollectiveMainloop = typename cutlass::gemm::collective::CollectiveBuilder< - cutlass::arch::Sm90, cutlass::arch::OpClassSparseTensorOp, - ElementAB, cutlass::layout::RowMajor, AlignmentAB, - ElementAB, cutlass::layout::ColumnMajor, AlignmentAB, + cutlass::arch::Sm90, cutlass::arch::OpClassSparseTensorOp, + ElementAB, cutlass::layout::RowMajor, AlignmentAB, + ElementAB, cutlass::layout::ColumnMajor, AlignmentAB, ElementAcc, TileShape, ClusterShape, Stages, KernelSchedule>::CollectiveOp;