Skip to content

[bugfix] cutlass_mla_decode is incorrectly declared when CUDA_VERSION… #9407

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
wants to merge 1 commit into
base: main
Choose a base branch
from
Open

[bugfix] cutlass_mla_decode is incorrectly declared when CUDA_VERSION… #9407

Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
28 changes: 11 additions & 17 deletions sgl-kernel/csrc/attention/cutlass_mla_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -28,22 +28,7 @@ limitations under the License.
#include "cutlass_sm100_mla/kernel/sm100_mla_tile_scheduler.hpp"

// clang-format off
#if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
void cutlass_mla_decode(
torch::Tensor const& out,
torch::Tensor const& q_nope,
torch::Tensor const& q_pe,
torch::Tensor const& kv_c_and_k_pe_cache,
torch::Tensor const& seq_lens,
torch::Tensor const& page_table,
torch::Tensor const& workspace,
int64_t num_kv_splits) {
TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
}
int64_t cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches, int64_t sm_count, int64_t num_kv_splits) {
TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
}
#else
#if defined(CUDA_VERSION) && CUDA_VERSION >= 12040
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

medium

To improve readability and maintainability, consider defining a macro for the CUDA version check at the top of the file. This avoids repeating the complex preprocessor condition throughout the file.

For example:

#if defined(CUDA_VERSION) && CUDA_VERSION >= 12040
#define SGL_HAS_CUTLASS_MLA_SUPPORT
#endif

You can then use this macro to simplify the conditional compilation blocks:

  1. This line becomes #ifdef SGL_HAS_CUTLASS_MLA_SUPPORT.
  2. The check inside cutlass_mla_decode becomes #ifndef SGL_HAS_CUTLASS_MLA_SUPPORT.
  3. The check inside cutlass_mla_get_workspace_size also becomes #ifndef SGL_HAS_CUTLASS_MLA_SUPPORT.


#define CUTLASS_CHECK(status) \
{ \
Expand Down Expand Up @@ -207,6 +192,8 @@ void runMla(
} \
}()

#endif

void cutlass_mla_decode(
torch::Tensor const& out,
torch::Tensor const& q_nope,
Expand All @@ -217,6 +204,9 @@ void cutlass_mla_decode(
torch::Tensor const& workspace,
double sm_scale,
int64_t num_kv_splits) {
#if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_decode");
#else
auto in_dtype = q_nope.dtype();
at::cuda::CUDAGuard device_guard{(char)q_nope.get_device()};
const cudaStream_t stream = at::cuda::getCurrentCUDAStream(q_nope.get_device());
Expand All @@ -243,9 +233,13 @@ void cutlass_mla_decode(
});
return true;
});
#endif
}

int64_t cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches, int64_t sm_count, int64_t num_kv_splits) {
#if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
TORCH_CHECK(false, "CUDA version must be >= 12.4 for cutlass_mla_get_workspace_size");
#else
// Workspace size depends on ElementAcc and ElementLSE (same as ElementAcc)
// which are float, so Element type here doesn't matter.
using MlaSm100Type = MlaSm100<cutlass::half_t, true>;
Expand All @@ -263,7 +257,7 @@ int64_t cutlass_mla_get_workspace_size(int64_t max_seq_len, int64_t num_batches,
MlaSm100Type::Fmha::set_split_kv(arguments);

return MlaSm100Type::Fmha::get_workspace_size(arguments);
#endif
}

#endif
// clang-format on