This is an automated email from the ASF dual-hosted git repository.
tlopex pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 865c2ea918 [Runtime] Fix CUDA build breaks in fp8 cutlass and thrust
(#19980)
865c2ea918 is described below
commit 865c2ea91887d4cf23b1161c15765c749061adea
Author: Ruihang Lai <[email protected]>
AuthorDate: Fri Jul 10 23:10:56 2026 -0400
[Runtime] Fix CUDA build breaks in fp8 cutlass and thrust (#19980)
The fp8 group-wise scaled GEMM kernels passed a braced DLDataType
initializer as the second argument of the two-argument TVM_FFI_ICHECK_EQ
macro (e.g. TVM_FFI_ICHECK_EQ(a->dtype, DLDataType{kDLFloat8_e4m3fn, 8,
1})). The preprocessor ignores brace grouping and splits on the commas
inside {...}, so it sees four arguments and fails to compile. Wrap the
initializer in parentheses so it is treated as a single macro argument.
thrust.cu calls args[i].cast<DLTensor*>() but did not include
<tvm/ffi/container/tensor.h>, which defines TypeTraits<DLTensor*>;
without it the cast fails template deduction. Add the include.
Both issues break the CUDA runtime build; with them fixed it compiles
cleanly with USE_CUTLASS and USE_THRUST enabled.
---
python/tvm/relax/backend/dispatch_sort_scan.py | 4 +---
.../contrib/cutlass/fp8_groupwise_scaled_gemm.cuh | 20 ++++++++++----------
.../cutlass/fp8_groupwise_scaled_group_gemm_sm100.cu | 12 ++++++------
src/runtime/extra/contrib/thrust/thrust.cu | 1 +
tests/python/relax/test_frontend_onnx.py | 8 ++++++--
5 files changed, 24 insertions(+), 21 deletions(-)
diff --git a/python/tvm/relax/backend/dispatch_sort_scan.py
b/python/tvm/relax/backend/dispatch_sort_scan.py
index 929e2e2781..76951718c4 100644
--- a/python/tvm/relax/backend/dispatch_sort_scan.py
+++ b/python/tvm/relax/backend/dispatch_sort_scan.py
@@ -86,9 +86,7 @@ class SortScanDispatcher(BackendDispatcher):
if self.is_gpu_target(tgt):
te_func = topi.gpu.searchsorted
out_dtype = "int32" if call.attrs.out_int32 else "int64"
- return self.builder_.call_te(
- te_func, boundaries, input_tensor, right, out_dtype
- )
+ return self.builder_.call_te(te_func, boundaries, input_tensor,
right, out_dtype)
if call.op.name == "relax.sort":
tgt = self._get_target(call.ty)
te_func = topi.sort
diff --git a/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_gemm.cuh
b/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_gemm.cuh
index 1af60af4da..540ef69999 100644
--- a/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_gemm.cuh
+++ b/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_gemm.cuh
@@ -66,11 +66,11 @@ void tvm_cutlass_fp8_groupwise_scaled_gemm_impl(Tensor a,
Tensor b, Tensor scale
TVM_FFI_ICHECK_EQ((n + block_size_0 - 1) / block_size_0, scales_b->shape[0]);
TVM_FFI_ICHECK_EQ(scales_b->shape[1] * block_size_1, k);
- TVM_FFI_ICHECK_EQ(a->dtype, DLDataType{kDLFloat8_e4m3fn, 8, 1});
- TVM_FFI_ICHECK_EQ(b->dtype, DLDataType{kDLFloat8_e4m3fn, 8, 1});
- TVM_FFI_ICHECK_EQ(scales_a->dtype, DLDataType{kDLFloat, 32, 1});
- TVM_FFI_ICHECK_EQ(scales_b->dtype, DLDataType{kDLFloat, 32, 1});
- TVM_FFI_ICHECK_EQ(workspace->dtype, DLDataType{kDLUInt, 8, 1});
+ TVM_FFI_ICHECK_EQ(a->dtype, (DLDataType{kDLFloat8_e4m3fn, 8, 1}));
+ TVM_FFI_ICHECK_EQ(b->dtype, (DLDataType{kDLFloat8_e4m3fn, 8, 1}));
+ TVM_FFI_ICHECK_EQ(scales_a->dtype, (DLDataType{kDLFloat, 32, 1}));
+ TVM_FFI_ICHECK_EQ(scales_b->dtype, (DLDataType{kDLFloat, 32, 1}));
+ TVM_FFI_ICHECK_EQ(workspace->dtype, (DLDataType{kDLUInt, 8, 1}));
int64_t workspace_nbytes =
workspace->shape[0] * ((workspace->dtype.bits * workspace->dtype.lanes +
7) / 8);
@@ -130,11 +130,11 @@ void tvm_cutlass_fp8_groupwise_scaled_bmm_impl(Tensor a,
Tensor b, Tensor scales
TVM_FFI_ICHECK_EQ(scales_b->shape[1] * block_size_0, n);
TVM_FFI_ICHECK_EQ(scales_b->shape[2] * block_size_1, k);
- TVM_FFI_ICHECK_EQ(a->dtype, DLDataType{kDLFloat8_e4m3fn, 8, 1});
- TVM_FFI_ICHECK_EQ(b->dtype, DLDataType{kDLFloat8_e4m3fn, 8, 1});
- TVM_FFI_ICHECK_EQ(scales_a->dtype, DLDataType{kDLFloat, 32, 1});
- TVM_FFI_ICHECK_EQ(scales_b->dtype, DLDataType{kDLFloat, 32, 1});
- TVM_FFI_ICHECK_EQ(workspace->dtype, DLDataType{kDLUInt, 8, 1});
+ TVM_FFI_ICHECK_EQ(a->dtype, (DLDataType{kDLFloat8_e4m3fn, 8, 1}));
+ TVM_FFI_ICHECK_EQ(b->dtype, (DLDataType{kDLFloat8_e4m3fn, 8, 1}));
+ TVM_FFI_ICHECK_EQ(scales_a->dtype, (DLDataType{kDLFloat, 32, 1}));
+ TVM_FFI_ICHECK_EQ(scales_b->dtype, (DLDataType{kDLFloat, 32, 1}));
+ TVM_FFI_ICHECK_EQ(workspace->dtype, (DLDataType{kDLUInt, 8, 1}));
int64_t workspace_nbytes =
workspace->shape[0] * ((workspace->dtype.bits * workspace->dtype.lanes +
7) / 8);
diff --git
a/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_group_gemm_sm100.cu
b/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_group_gemm_sm100.cu
index 6bd9f45ab2..d37bc35568 100644
--- a/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_group_gemm_sm100.cu
+++ b/src/runtime/extra/contrib/cutlass/fp8_groupwise_scaled_group_gemm_sm100.cu
@@ -57,12 +57,12 @@ void tvm_fp8_groupwise_scaled_group_gemm_sm100(Tensor a,
Tensor b, Tensor scales
TVM_FFI_ICHECK_EQ((n + block_size_0 - 1) / block_size_0, scales_b->shape[1]);
TVM_FFI_ICHECK_EQ((k + block_size_1 - 1) / block_size_1, scales_b->shape[2]);
- TVM_FFI_ICHECK_EQ(a->dtype, DLDataType{kDLFloat8_e4m3fn, 8, 1});
- TVM_FFI_ICHECK_EQ(b->dtype, DLDataType{kDLFloat8_e4m3fn, 8, 1});
- TVM_FFI_ICHECK_EQ(scales_a->dtype, DLDataType{kDLFloat, 32, 1});
- TVM_FFI_ICHECK_EQ(scales_b->dtype, DLDataType{kDLFloat, 32, 1});
- TVM_FFI_ICHECK_EQ(indptr->dtype, DLDataType{kDLInt, 64, 1});
- TVM_FFI_ICHECK_EQ(workspace->dtype, DLDataType{kDLUInt, 8, 1});
+ TVM_FFI_ICHECK_EQ(a->dtype, (DLDataType{kDLFloat8_e4m3fn, 8, 1}));
+ TVM_FFI_ICHECK_EQ(b->dtype, (DLDataType{kDLFloat8_e4m3fn, 8, 1}));
+ TVM_FFI_ICHECK_EQ(scales_a->dtype, (DLDataType{kDLFloat, 32, 1}));
+ TVM_FFI_ICHECK_EQ(scales_b->dtype, (DLDataType{kDLFloat, 32, 1}));
+ TVM_FFI_ICHECK_EQ(indptr->dtype, (DLDataType{kDLInt, 64, 1}));
+ TVM_FFI_ICHECK_EQ(workspace->dtype, (DLDataType{kDLUInt, 8, 1}));
if (out->dtype == DLDataType{kDLFloat, 16, 1}) {
using Dtype = cutlass::half_t;
diff --git a/src/runtime/extra/contrib/thrust/thrust.cu
b/src/runtime/extra/contrib/thrust/thrust.cu
index f534e1fc76..cc530a8612 100644
--- a/src/runtime/extra/contrib/thrust/thrust.cu
+++ b/src/runtime/extra/contrib/thrust/thrust.cu
@@ -31,6 +31,7 @@
#include <thrust/scan.h>
#include <thrust/sequence.h>
#include <thrust/sort.h>
+#include <tvm/ffi/container/tensor.h>
#include <tvm/ffi/dtype.h>
#include <tvm/ffi/extra/c_env_api.h>
#include <tvm/ffi/extra/cuda/base.h>
diff --git a/tests/python/relax/test_frontend_onnx.py
b/tests/python/relax/test_frontend_onnx.py
index f96984e423..a076f01c4a 100644
--- a/tests/python/relax/test_frontend_onnx.py
+++ b/tests/python/relax/test_frontend_onnx.py
@@ -6563,7 +6563,9 @@ def test_pad(dynamic):
if axes is not None:
axes = np.array(axes, dtype=np.int64)
node_inputs = ["input", "pads", "", "axes"]
- initializer.append(helper.make_tensor("axes",
TensorProto.INT64, (len(axes),), axes))
+ initializer.append(
+ helper.make_tensor("axes", TensorProto.INT64,
(len(axes),), axes)
+ )
node = helper.make_node("Pad", inputs=node_inputs,
outputs=["output"], mode=mode)
graph = helper.make_graph(
@@ -6628,7 +6630,9 @@ def test_pad(dynamic):
verify_pad(
input_shape,
pads,
- _make_pad_expected_ir(input_shape, pads, mode=mode, value=value,
opset=opset, axes=axes),
+ _make_pad_expected_ir(
+ input_shape, pads, mode=mode, value=value, opset=opset,
axes=axes
+ ),
mode,
value,
opset,
