Skip to content

feat: add SM110 support #13

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
AlpinDale wants to merge 1 commit into
base: main
Choose a base branch
from
+23 −8
Open

feat: add SM110 support #13

Show file tree
Hide file tree
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
4 changes: 2 additions & 2 deletions qutlass/csrc/bindings.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -308,7 +308,7 @@ std::tuple<torch::Tensor, torch::Tensor> fusedQuantizeMxAbsMax(torch::Tensor con
} else if(HAD_GS==64){
fusedQuantizeMxAbsMaxHad64_host(OUT, OUT_sf, A, B);
} else if(HAD_GS==128){
#if TARGET_CUDA_ARCH == 100
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110
auto opts = torch::TensorOptions().dtype(torch::kFloat).device(A.device());
auto global_scale = torch::tensor(0.0f, opts); //FIXME: add input global_scale to interface for consistency
fusedQuantizeMxAbsMax_host_sm100(OUT, OUT_sf, A, B, global_scale);
Expand Down Expand Up @@ -399,7 +399,7 @@ std::tuple<torch::Tensor, torch::Tensor> fusedQuantizeNvAbsMax(torch::Tensor con
} else if(HAD_GS==64){
fusedQuantizeNvAbsMaxHad64_host(OUT, OUT_sf, A, B, global_scale);
} else if(HAD_GS==128){
#if TARGET_CUDA_ARCH == 100
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110
fusedQuantizeNvAbsMax_host_sm100(OUT, OUT_sf, A, B, global_scale);
#elif TARGET_CUDA_ARCH == 120
fusedQuantizeNvAbsMaxHad128_host(OUT, OUT_sf, A, B, global_scale);
Expand Down
2 changes: 1 addition & 1 deletion qutlass/csrc/fused_quantize_mx_sm100.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -205,7 +205,7 @@ void fusedQuantizeMxAbsMax_host_sm100(torch::Tensor& D,
torch::Tensor const& B,
torch::Tensor const& global_scale)
{
#if TARGET_CUDA_ARCH == 100
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110
int32_t M = A.numel() / 128;
int32_t N = B.size(1);
int32_t K = 128;
Expand Down
2 changes: 1 addition & 1 deletion qutlass/csrc/fused_quantize_nv_sm100.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -205,7 +205,7 @@ void fusedQuantizeNvAbsMax_host_sm100(torch::Tensor& D,
torch::Tensor const& B,
torch::Tensor const& global_scale)
{
#if TARGET_CUDA_ARCH == 100
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110
int32_t M = A.numel() / 128;
int32_t N = B.size(1);
int32_t K = 128;
Expand Down
8 changes: 4 additions & 4 deletions qutlass/csrc/gemm.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -200,7 +200,7 @@ void matmul_host_mxf4_bf16_tn(torch::Tensor& D,
using LayoutBTag = cutlass::layout::ColumnMajor;
static constexpr int AlignmentB = 128;

#if TARGET_CUDA_ARCH == 100 //TODO: improve tuning
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110 //TODO: improve tuning
using ArchTag = cutlass::arch::Sm100;
if(m<=16){
using MmaTileShape = Shape<_128,_128,_256>;
Expand Down Expand Up @@ -276,7 +276,7 @@ void matmul_host_nvf4_bf16_tn(torch::Tensor& D,
using LayoutBTag = cutlass::layout::ColumnMajor;
static constexpr int AlignmentB = 32;

#if TARGET_CUDA_ARCH == 100 //TODO: improve tuning
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110 //TODO: improve tuning
using ArchTag = cutlass::arch::Sm100;
if(m<=16){
using MmaTileShape = Shape<_128,_128,_256>;
Expand Down Expand Up @@ -354,7 +354,7 @@ void matmul_host_mxf8_bf16_tn(torch::Tensor& D,
using LayoutBTag = cutlass::layout::ColumnMajor;
static constexpr int AlignmentB = 16;

#if TARGET_CUDA_ARCH == 100
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110
using ArchTag = cutlass::arch::Sm100;

if(m<=8192){
Expand Down Expand Up @@ -414,7 +414,7 @@ void matmul_host_mxf8_bf16_nn(torch::Tensor& D,
using LayoutBTag = cutlass::layout::ColumnMajor;
static constexpr int AlignmentB = 16;

#if TARGET_CUDA_ARCH == 100
#if TARGET_CUDA_ARCH == 100 || TARGET_CUDA_ARCH == 101 || TARGET_CUDA_ARCH == 110
using ArchTag = cutlass::arch::Sm100;

if(m<=8192){
Expand Down
15 changes: 15 additions & 0 deletions setup.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -57,6 +57,21 @@ def get_cuda_arch_flags():
"arch=compute_120a,code=sm_120a",
"-gencode",
"arch=compute_100a,code=sm_100a",
]
# Thor / GB10 (SM 10.1, renamed SM 11.0 in CUDA 13.0) shares the SM100
# tcgen05 block-scaled fp4 path. Only relevant when building on such a
# device (this is a per-machine build), so gate on the local capability
# rather than the CPU arch -- GB200/GH200 are also aarch64 but not Thor.
# Emit its SASS under whichever name the active toolkit uses: sm_110a on
# CUDA >= 13.0, sm_101a on <= 12.9.
cuda_ver = torch.version.cuda
if cc in (101, 110) and cuda_ver is not None:
cuda_major, cuda_minor = (int(v) for v in cuda_ver.split(".")[:2])
if (cuda_major, cuda_minor) >= (13, 0):
flags += ["-gencode", "arch=compute_110a,code=sm_110a"]
Comment thread
AlpinDale marked this conversation as resolved.
else:
flags += ["-gencode", "arch=compute_101a,code=sm_101a"]
flags += [
"--expt-relaxed-constexpr",
"--use_fast_math",
"-std=c++17",
Expand Down