AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Support passing macros to hipLaunchKernelGGL

Browse Source 
This makes hipLaunchKernelGGL take a variable argument list, that will be
expanded before being fed to hipLaunchKernelGGLInternal.

This is different from 4dfef9a25b.

We try to accomodate the case when a kernel template has multiple
type parameters.

Change-Id: I87577d402c92b0f3b51e298f8293f4065e1f6de8


[ROCm/clr commit: 133251a6e1]
This commit is contained in:
Ronak Chauhan
2020-06-25 07:23:28 -04:00
committed by Aaron En Ye Shi
parent a2e2f9b8c6
commit 845c0a5b8a
3 changed files with 24 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/include/hip/hcc_detail/hip_runtime.h
+4 -2
View File
@@ -435,10 +435,12 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
hipLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream);
}
#else
#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \
do { \
kernelName<<<(numblocks), (numthreads), (memperblock), (streamId)>>>(__VA_ARGS__); \
kernelName<<<(numBlocks), (numThreads), (memPerBlock), (streamId)>>>(__VA_ARGS__); \
} while (0)
#define hipLaunchKernelGGL(kernelName, ...) hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__)
#endif
#include <hip/hip_runtime_api.h>
projects/clr/hipamd/include/hip/nvcc_detail/hip_runtime.h
+4 -2
View File
@@ -31,11 +31,13 @@ THE SOFTWARE.
typedef int hipLaunchParm;
#define hipLaunchKernelGGL(kernelName, numblocks, numthreads, memperblock, streamId, ...) \
#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \
do { \
kernelName<<<numblocks, numthreads, memperblock, streamId>>>(__VA_ARGS__); \
kernelName<<<numBlocks, numThreads, memPerBlock, streamId>>>(__VA_ARGS__); \
} while (0)
#define hipLaunchKernelGGL(kernelName, ...) hipLaunchKernelGGLInternal((kernelName), __VA_ARGS__)
#define hipReadModeElementType cudaReadModeElementType
#ifdef __CUDA_ARCH__