This is an automated email from the ASF dual-hosted git repository.
junrushao pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm-ffi.git
The following commit(s) were added to refs/heads/main by this push:
new a23c5a0 [Minor] Fix missing function ref in type traits (#186)
a23c5a0 is described below
commit a23c5a0350fcab899d1623b4db7fa3b84f12daaa
Author: DarkSharpness <[email protected]>
AuthorDate: Wed Oct 22 05:32:21 2025 +0800
[Minor] Fix missing function ref in type traits (#186)
Before this fix, here's one failure case (modified from `add_one_cuda
example, and actually I encounter this error in a similar real-world
case.):
```cpp
// File: compile/add_one_cuda.cu
#include <tvm/ffi/container/tensor.h>
#include <tvm/ffi/extra/c_env_api.h>
#include <tvm/ffi/function.h>
namespace tvm_ffi_example_cuda {
__global__ void AddOneKernel(float* x, float* y, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
y[idx] = x[idx] + 1;
}
}
void AddOne(tvm::ffi::TensorView x, tvm::ffi::TensorView y) {
int64_t n = x.size(0);
float* x_data = static_cast<float*>(x.data_ptr());
float* y_data = static_cast<float*>(y.data_ptr());
int64_t threads = 256;
int64_t blocks = (n + threads - 1) / threads;
cudaStream_t stream =
static_cast<cudaStream_t>(TVMFFIEnvGetStream(x.device().device_type,
x.device().device_id));
AddOneKernel<<<blocks, threads, 0, stream>>>(x_data, y_data, n);
}
// this will crash, because extra () will decltype as a reference type
TVM_FFI_DLL_EXPORT_TYPED_FUNC(add_one_cuda, (tvm_ffi_example_cuda::AddOne));
// this will also crash
static constexpr auto& func = tvm_ffi_example_cuda::AddOne;
TVM_FFI_DLL_EXPORT_TYPED_FUNC(add_one_cuda_1, func);
} // namespace tvm_ffi_example_cuda
```
---
include/tvm/ffi/function_details.h | 2 ++
1 file changed, 2 insertions(+)
diff --git a/include/tvm/ffi/function_details.h
b/include/tvm/ffi/function_details.h
index e7766fd..e7db063 100644
--- a/include/tvm/ffi/function_details.h
+++ b/include/tvm/ffi/function_details.h
@@ -113,6 +113,8 @@ template <typename R, typename... Args>
struct FunctionInfo<R(Args...), void> : FuncFunctorImpl<R, Args...> {};
template <typename R, typename... Args>
struct FunctionInfo<R (*)(Args...), void> : FuncFunctorImpl<R, Args...> {};
+template <typename R, typename... Args>
+struct FunctionInfo<R (&)(Args...), void> : FuncFunctorImpl<R, Args...> {};
// Support pointer-to-member functions used in reflection (e.g. &Class::method)
template <typename Class, typename R, typename... Args>
struct FunctionInfo<R (Class::*)(Args...),
std::enable_if_t<std::is_base_of_v<Object, Class>>>
