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.

Change-Id: Id76e2bf91acd5d68f56a24fc39f219f2eeb06d33
This commit is contained in:
Ronak Chauhan
2020-06-16 02:49:12 -04:00
committed by Ronak Nilesh Chauhan
parent 63051ca2e1
commit 961717879d
3 changed files with 25 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
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(...) hipLaunchKernelGGLInternal(__VA_ARGS__)
#endif
#include <hip/hip_runtime_api.h>
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(...) hipLaunchKernelGGLInternal(__VA_ARGS__)
#define hipReadModeElementType cudaReadModeElementType
#ifdef __CUDA_ARCH__