Skip to content

[Build] Fix cuda12.8 build error in nvfp4_scaled_mm_kernels.cu #4953

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
Mar 31, 2025
Merged

[Build] Fix cuda12.8 build error in nvfp4_scaled_mm_kernels.cu #4953

Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
e1d99c0
Disable build sm_100 while using cuda12.8
guoyuhong Mar 31, 2025
d009eba
Try to fix the build error
guoyuhong Mar 31, 2025
ce4c2dd
Fix typo
guoyuhong Mar 31, 2025
7bf058c
Revert unnecessary code change
guoyuhong Mar 31, 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
6 changes: 3 additions & 3 deletions sgl-kernel/csrc/gemm/nvfp4_scaled_mm_kernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -159,7 +159,7 @@ typename T::Gemm::Arguments args_from_options(
using StrideA = typename T::StrideA;
using StrideB = typename T::StrideB;
using StrideD = typename T::StrideD;
using Sm100BlkScaledConfig = typename T::Gemm::GemmKernel::CollectiveMainloop::Sm100BlkScaledConfig;
using Sm1xxBlkScaledConfig = typename T::Gemm::GemmKernel::CollectiveMainloop::Sm1xxBlkScaledConfig;

int m = static_cast<int>(M);
int n = static_cast<int>(N);
Expand All @@ -168,8 +168,8 @@ typename T::Gemm::Arguments args_from_options(
auto stride_B = cutlass::make_cute_packed_stride(StrideB{}, {n, k, 1});
auto stride_D = cutlass::make_cute_packed_stride(StrideD{}, {m, n, 1});

auto layout_SFA = Sm100BlkScaledConfig::tile_atom_to_shape_SFA(cute::make_shape(m, n, k, 1));
auto layout_SFB = Sm100BlkScaledConfig::tile_atom_to_shape_SFB(cute::make_shape(m, n, k, 1));
auto layout_SFA = Sm1xxBlkScaledConfig::tile_atom_to_shape_SFA(cute::make_shape(m, n, k, 1));
auto layout_SFB = Sm1xxBlkScaledConfig::tile_atom_to_shape_SFB(cute::make_shape(m, n, k, 1));

typename T::Gemm::Arguments arguments{
cutlass::gemm::GemmUniversalMode::kGemm,
Expand Down
Loading