AI/rocm-systems
2
0
Fork 0
Koodi Ongelmat Pull-pyynnöt Actions Paketit Projektit Julkaisut Wiki Toiminta

Revert "Support passing macros to hipLaunchKernelGGL"

Selaa lähdekoodia 
This reverts commit b2edee4693.

Reason for revert: This patch breaks ROCPrim tests

Change-Id: Ib2235f719861c9f4317c33e86b6c1f8bc669cfd4


[ROCm/hip commit: b7101af203]
This commit is contained in:
Ronak Nilesh Chauhan
2020-06-24 02:38:00 -04:00
vanhempi bab86e0e28
commit b14dacc35f
3 muutettua tiedostoa jossa 4 lisäystä ja 25 poistoa
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/hip_runtime.h
+2 -4
Näytä tiedosto
@@ -435,12 +435,10 @@ void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
hipLaunchKernel(k, numBlocks, dimBlocks, _Args, sharedMemBytes, stream);
}
#else
#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \
#define hipLaunchKernelGGL(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>
projects/hip/include/hip/nvcc_detail/hip_runtime.h
+2 -4
Näytä tiedosto
@@ -31,13 +31,11 @@ THE SOFTWARE.
typedef int hipLaunchParm;
#define hipLaunchKernelGGLInternal(kernelName, numBlocks, numThreads, memPerBlock, streamId, ...) \
#define hipLaunchKernelGGL(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__
projects/hip/tests/src/kernel/hipLaunchParm.cpp
-17
Näytä tiedosto
@@ -571,10 +571,6 @@ __global__ void hipLaunchKernelStructFunc21(
__global__ void vAdd(float* a) {}
template<class T1, class T2>
__global__ void myKernel(T1 a, T2 b) {}
//---
// Some wrapper macro for testing:
#define WRAP(...) __VA_ARGS__
@@ -917,19 +913,6 @@ int main() {
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), 1, 0, 0, Ad);
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), dim3(1024), dim3(1), 0, 0, Ad);
// Test: Passing macro to hipLaunchKernelGGL
#define KERNEL_CONFIG dim3(1024), dim3(1), 0, 0
hipLaunchKernelGGL(HIP_KERNEL_NAME(vAdd), KERNEL_CONFIG, Ad);
// Test: Same thing with templates:
// Additional parenthesis required case of templates with multiple type paramters.
int a;
float b;
hipLaunchKernelGGL(HIP_KERNEL_NAME((myKernel<int, float>)), KERNEL_CONFIG, a, b);
#define TYPE_PARAM_CONFIG int, float
hipLaunchKernelGGL(HIP_KERNEL_NAME((myKernel<TYPE_PARAM_CONFIG>)), KERNEL_CONFIG, a, b);
// Test: Passing hipLaunchKernelGGL inside another macro:
float e0;
MY_LAUNCH_MACRO(hipLaunchKernelGGL(vAdd, dim3(1024),