From eb9c852dc57f6a41059eb2013b8cebb5a59db3d6 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Mon, 17 Sep 2018 11:19:35 -0400 Subject: [PATCH] Fix hipLaunchKernelGGL for hip-clang Do not decay function pointer type of the kernel argument passed to hipLaunchKernelGGL and hipLaunchKernel, otherwise some type information is lost which may cause type inference failure for the template. This issue caused compilation error of FeatureLPPooling in Caffe2/PyTorch and this patch fixes that. [ROCm/clr commit: 0471ce2da65b48fe4504906f556ec86bc3ce0d63] --- projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h index e5f0fb52fa..d8093e6646 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h @@ -334,13 +334,13 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri typedef int hipLaunchParm; template -inline void hipLaunchKernelGGL(F kernelName, const dim3& numblocks, const dim3& numthreads, +inline void hipLaunchKernelGGL(F&& kernelName, const dim3& numblocks, const dim3& numthreads, unsigned memperblock, hipStream_t streamId, Args... args) { kernelName<<>>(args...); } template -inline void hipLaunchKernel(F kernel, const dim3& numBlocks, const dim3& dimBlocks, +inline void hipLaunchKernel(F&& kernel, const dim3& numBlocks, const dim3& dimBlocks, std::uint32_t groupMemBytes, hipStream_t stream, Args... args) { hipLaunchKernelGGL(kernel, numBlocks, dimBlocks, groupMemBytes, stream, hipLaunchParm{}, std::move(args)...);