From 9fe28c2987348e481f1e3f5d53b640e92f3b8cac Mon Sep 17 00:00:00 2001 From: Kokoro2336 <2529677678@qq.com> Date: Thu, 23 Apr 2026 19:03:34 +0800 Subject: [PATCH 1/5] fix: missing symbol due to version mismatch. --- src/runtime/contrib/cublas/cublas.cc | 8 +++++++- 1 file changed, 7 insertions(+), 1 deletion(-) diff --git a/src/runtime/contrib/cublas/cublas.cc b/src/runtime/contrib/cublas/cublas.cc index 6e667ee37878..b02c4748e06b 100644 --- a/src/runtime/contrib/cublas/cublas.cc +++ b/src/runtime/contrib/cublas/cublas.cc @@ -169,10 +169,14 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, ab_type = CUDA_R_16BF; } else if (TypeMatch(A->dtype, kDLInt, 8)) { ab_type = CUDA_R_8I; - } else if (TypeMatch(A->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)) { + } + +#if CUDART_VERSION >= 11800 + else if (TypeMatch(A->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)) { TVM_FFI_ICHECK(TypeMatch(B->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)); ab_type = CUDA_R_8F_E4M3; } +#endif if (TypeMatch(C->dtype, kDLFloat, 16)) { c_type = CUDA_R_16F; @@ -201,6 +205,7 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, &bias->data, sizeof(float*))); } +#if CUDART_VERSION >= 11800 if (scaleA != nullptr) { auto scaleA_data = static_cast(scaleA->data) + scaleA->byte_offset; CHECK_CUBLAS_ERROR(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_A_SCALE_POINTER, @@ -211,6 +216,7 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, CHECK_CUBLAS_ERROR(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_B_SCALE_POINTER, &scaleB_data, sizeof(float*))); } +#endif if (epilogue != CUBLASLT_EPILOGUE_DEFAULT) { CHECK_CUBLAS_ERROR(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_EPILOGUE, From da1bc1afde18d2c99ee041922f41bdb6d0b17ec4 Mon Sep 17 00:00:00 2001 From: Kokoro2336 <2529677678@qq.com> Date: Thu, 23 Apr 2026 19:40:55 +0800 Subject: [PATCH 2/5] fix: error handling of unsupported version. --- src/runtime/contrib/cublas/cublas.cc | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/src/runtime/contrib/cublas/cublas.cc b/src/runtime/contrib/cublas/cublas.cc index b02c4748e06b..d26c6ce57944 100644 --- a/src/runtime/contrib/cublas/cublas.cc +++ b/src/runtime/contrib/cublas/cublas.cc @@ -169,14 +169,14 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, ab_type = CUDA_R_16BF; } else if (TypeMatch(A->dtype, kDLInt, 8)) { ab_type = CUDA_R_8I; - } - + } else if (TypeMatch(A->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)) { #if CUDART_VERSION >= 11800 - else if (TypeMatch(A->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)) { TVM_FFI_ICHECK(TypeMatch(B->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)); ab_type = CUDA_R_8F_E4M3; - } +#else + TVM_FFI_THROW(InternalError) << "Float8 (E4M3) is only supported in CUDA 11.8 and above."; #endif + } if (TypeMatch(C->dtype, kDLFloat, 16)) { c_type = CUDA_R_16F; From cbeb4cb980099a98d567f445214b309d9e727244 Mon Sep 17 00:00:00 2001 From: Kokoro2336 <2529677678@qq.com> Date: Thu, 23 Apr 2026 19:50:46 +0800 Subject: [PATCH 3/5] ci: retrigger ci. From e8f7d6b3cd7c032c4921aa6ecf1cf29fdfb71726 Mon Sep 17 00:00:00 2001 From: Kokoro2336 <2529677678@qq.com> Date: Thu, 23 Apr 2026 19:52:52 +0800 Subject: [PATCH 4/5] fix: error handling for scale pointers. --- src/runtime/contrib/cublas/cublas.cc | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/src/runtime/contrib/cublas/cublas.cc b/src/runtime/contrib/cublas/cublas.cc index d26c6ce57944..c122797c29ec 100644 --- a/src/runtime/contrib/cublas/cublas.cc +++ b/src/runtime/contrib/cublas/cublas.cc @@ -216,6 +216,10 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, CHECK_CUBLAS_ERROR(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_B_SCALE_POINTER, &scaleB_data, sizeof(float*))); } +#else + if (scaleA != nullptr || scaleB != nullptr) { + TVM_FFI_THROW(InternalError) << "Scaling pointers are only supported in CUDA 11.8 and above."; + } #endif if (epilogue != CUBLASLT_EPILOGUE_DEFAULT) { From 067a2c72a314d9c05f6c44d080c0becbd855fe06 Mon Sep 17 00:00:00 2001 From: Kokoro2336 <2529677678@qq.com> Date: Thu, 23 Apr 2026 19:58:04 +0800 Subject: [PATCH 5/5] fix: version number. --- src/runtime/contrib/cublas/cublas.cc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/src/runtime/contrib/cublas/cublas.cc b/src/runtime/contrib/cublas/cublas.cc index c122797c29ec..dcaf93d2da2e 100644 --- a/src/runtime/contrib/cublas/cublas.cc +++ b/src/runtime/contrib/cublas/cublas.cc @@ -170,7 +170,7 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, } else if (TypeMatch(A->dtype, kDLInt, 8)) { ab_type = CUDA_R_8I; } else if (TypeMatch(A->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)) { -#if CUDART_VERSION >= 11800 +#if CUDART_VERSION >= 11080 TVM_FFI_ICHECK(TypeMatch(B->dtype, DataType::TypeCode::kFloat8_e4m3fn, 8)); ab_type = CUDA_R_8F_E4M3; #else @@ -205,7 +205,7 @@ void CallCublasLt(cublasLtHandle_t hdl, cudaStream_t stream, &bias->data, sizeof(float*))); } -#if CUDART_VERSION >= 11800 +#if CUDART_VERSION >= 11080 if (scaleA != nullptr) { auto scaleA_data = static_cast(scaleA->data) + scaleA->byte_offset; CHECK_CUBLAS_ERROR(cublasLtMatmulDescSetAttribute(op_desc, CUBLASLT_MATMUL_DESC_A_SCALE_POINTER,