Skip to content

Disable kernel cutlass_mla_decode on SM103 #10058

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

Merged
merged 4 commits into from
Sep 6, 2025
+8 −2
Merged

Disable kernel cutlass_mla_decode on SM103 #10058

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
7911194
Disable kernel cutlass_mla_decode on SM103
hlu1 Sep 5, 2025
0c5ba07
Merge branch 'main' into mla
zhyncs Sep 6, 2025
8ca0297
Merge branch 'main' into mla
zhyncs Sep 6, 2025
30cecb4
Merge branch 'main' into mla
zhyncs Sep 6, 2025
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
5 changes: 5 additions & 0 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 @@ -26,6 +26,7 @@ limitations under the License.

#include "cutlass_sm100_mla/device/sm100_mla.hpp"
#include "cutlass_sm100_mla/kernel/sm100_mla_tile_scheduler.hpp"
#include "utils.h"

// clang-format off
#if !defined(CUDA_VERSION) || CUDA_VERSION < 12040
Expand Down Expand Up @@ -217,6 +218,10 @@ void cutlass_mla_decode(
torch::Tensor const& workspace,
double sm_scale,
int64_t num_kv_splits) {
auto sm_version = getSMVersion();
// On SM103a, half of the accuracy tests are failing.
TORCH_CHECK(sm_version == 100, "cutlass_mla_decode is only supported on compute capability 10.0, but found sm version ", sm_version);

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 Down
5 changes: 3 additions & 2 deletions sgl-kernel/tests/test_cutlass_mla.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,9 +4,10 @@
from sgl_kernel import cutlass_mla_decode, cutlass_mla_get_workspace_size
from torch import Tensor

if torch.cuda.get_device_capability() < (10, 0):
# Disable tests on SM103 until the accuracy issues are fixed.
if torch.cuda.get_device_capability() != (10, 0):
pytest.skip(
reason="Cutlass MLA Requires compute capability of 10 or above.",
reason="Cutlass MLA Requires compute capability of 10.",
allow_module_level=True,
)

Expand Down
Loading