DarkSharpness opened a new pull request, #186:
URL: https://github.com/apache/tvm-ffi/pull/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, func);
} // namespace tvm_ffi_example_cuda
```
--
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
To unsubscribe, e-mail: [email protected]
For queries about this service, please contact Infrastructure at:
[email protected]
---------------------------------------------------------------------
To unsubscribe, e-mail: [email protected]
For additional commands, e-mail: [email protected]
