I was trying to port a CUDA project into ROCm platform which make use of C++ templates extensively. During this process, I am getting following compilation error
/root/warp-ctc/include/detail/gpu_ctc.h:381:5: error: no matching function for call to 'hipLaunchKernelGGL'
hipLaunchKernelGGL((prepare_stable_SM_kernel<ProbT, VT>), dim3(grid_size), dim3(NT), 0, stream_, ctc_helper::identity<ProbT>(), probs_,
^~~~~~~~~~~~~~~~~~
.....
.....
/opt/rocm/hip/include/hip/hcc_detail/functional_grid_launch.hpp:138:13: note: candidate function [with Args = <ctc_helper::identity<float, float>, float *,
float *, int, int>, F = void (*)(ctc_helper::identity<float, float>, float *, float *, int, int)] not viable: no overload of 'prepare_stable_SM_kernel'
matching 'void (*)(ctc_helper::identity<float, float>, float *, float *, int, int)' for 1st argument
inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
for following piece of C++ code
hipLaunchKernelGGL((prepare_stable_SM_kernel<ProbT, VT>),
dim3(grid_size),
dim3(NT),
0,
stream_,
ctc_helper::identity<ProbT>(),
probs_,
denoms_,
out_dim_,
num_elements);
where hipLaunchKernelGGL is defined as
template <typename... Args, typename F = void (*)(Args...)>
inline void hipLaunchKernelGGL(F kernel,
const dim3& numBlocks,
const dim3& dimBlocks,
std::uint32_t sharedMemBytes,
hipStream_t stream,
Args... args) {
// ...
// ...
}
and prepare_stable_SM_kernel is defined as
template <typename ProbT, int VT = 1, typename Op>
__global__ void prepare_stable_SM_kernel(Op f, ProbT* probs,
const ProbT* const col_max,
int alphabet_size,
int count) {
// ...
}
Anyone please help me by providing some hints to fix this problem.
This error was fixed by adding a third parameter ctc_helper::identity<ProbT> to prepare_stable_SM_kernel function