From 0471ce2da65b48fe4504906f556ec86bc3ce0d63 Mon Sep 17 00:00:00 2001
From: Yaxun Sam Liu <yaxun.liu@amd.com>
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.
---
 hipamd/include/hip/hcc_detail/hip_runtime.h | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_runtime.h b/hipamd/include/hip/hcc_detail/hip_runtime.h
index e5f0fb52fa..d8093e6646 100644
--- a/hipamd/include/hip/hcc_detail/hip_runtime.h
+++ b/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 <typename... Args, typename F = void (*)(Args...)>
-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<<<numblocks, numthreads, memperblock, streamId>>>(args...);
 }
 
 template <typename... Args, typename F = void (*)(hipLaunchParm, Args...)>
-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)...);