| [general] |
| name = "fp8-gemm" |
| version = 1 |
| edition = 5 |
| license = "Apache-2.0" |
| backends = ["cuda"] |
|
|
| [general.cuda] |
| minver = "12.8" |
|
|
| [general.hub] |
| repo-id = "flashrt/fp8-gemm" |
|
|
| [torch] |
| include = ["csrc"] |
| src = [ |
| "torch-ext/torch_binding.cpp", |
| "torch-ext/torch_binding.h", |
| ] |
|
|
| [kernel.fp8_gemm_sm110] |
| backend = "cuda" |
| cuda-capabilities = ["11.0a"] |
| cuda-flags = [ |
| "--expt-relaxed-constexpr", |
| "-O3", |
| "--use_fast_math", |
| ] |
| cuda-minver = "13" |
| depends = [ |
| "torch", |
| "cutlass_4_5", |
| ] |
| include = ["csrc"] |
| src = [ |
| "csrc/gemm_types_sm110.h", |
| "csrc/cutlass_sm110_fp8_gemm.cu", |
| "csrc/cutlass_sm110_fp8_gemm.cuh", |
| "csrc/cublaslt_fp8_bias_sm110.cu", |
| "csrc/cublaslt_fp8_bias_sm110.cuh", |
| ] |
|
|
| [kernel.fp8_gemm_sm89] |
| backend = "cuda" |
| cuda-capabilities = ["8.9"] |
| cuda-minver = "12.8" |
| depends = ["torch"] |
| include = ["csrc"] |
| src = [ |
| "csrc/fp8_block128_gemm_mma_sm89.cu", |
| "csrc/fp8_block128_gemm_mma_sm89.cuh", |
| "csrc/fp8_bs_gemm_device.cuh", |
| "csrc/fp8_gemv_m1_sm89.cu", |
| "csrc/fp8_gemv_m1_sm89.cuh", |
| ] |
|
|
| [kernel.fp8_gemm] |
| backend = "cuda" |
| cuda-capabilities = ["12.0a"] |
| cuda-minver = "12.8" |
| depends = [ |
| "torch", |
| "cutlass_4_0", |
| ] |
| include = ["csrc"] |
| src = [ |
| "csrc/fp8_gemv_m1_sm120.cu", |
| "csrc/fp8_gemv_m1_sm120.cuh", |
| "csrc/fp8_smallM_handtuned_sm120.cu", |
| "csrc/fp8_smallM_handtuned_sm120.cuh", |
| "csrc/fp8_smallM_handtuned_ldmatrix_sm120.cu", |
| "csrc/fp8_smallM_handtuned_ldmatrix_sm120.cuh", |
| "csrc/cutlass_sm120_block128_fp8_gemm.cu", |
| "csrc/cutlass_sm120_block128_fp8_gemm.cuh", |
| ] |
|
|
