File size: 1,549 Bytes
9022070 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 | [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",
]
|