From b14dacc35f0fca0040f7c5275e18e175e70ceea3 Mon Sep 17 00:00:00 2001
From: Ronak Nilesh Chauhan <RonakNilesh.Chauhan@amd.com>
Date: Wed, 24 Jun 2020 02:38:00 -0400
Subject: [PATCH] Revert "Support passing macros to hipLaunchKernelGGL"

This reverts commit b2edee4693113411e4e0207374048cebf6809547.

Reason for revert: This patch breaks ROCPrim tests

Change-Id: Ib2235f719861c9f4317c33e86b6c1f8bc669cfd4


[ROCm/hip commit: b7101af203d3288b323a4ad1b29392e4e3b6c778]
---
 .../hip/include/hip/hcc_detail/hip_runtime.h    |  6 ++----
 .../hip/include/hip/nvcc_detail/hip_runtime.h   |  6 ++----
 projects/hip/tests/src/kernel/hipLaunchParm.cpp | 17 -----------------
 3 files changed, 4 insertions(+), 25 deletions(-)

diff --git a/projects/hip/include/hip/hcc_detail/hip_runtime.h b/projects/hip/include/hip/hcc_detail/hip_runtime.h
index c35a435fea..6c7850fbbc 100644
--- a/projects/hip/include/hip/hcc_detail/hip_runtime.h
+++ b/projects/hip/include/hip/hcc_detail/hip_runtime.h
@@ -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>
diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime.h b/projects/hip/include/hip/nvcc_detail/hip_runtime.h
index 8443e3eb2a..d159e116c0 100644
--- a/projects/hip/include/hip/nvcc_detail/hip_runtime.h
+++ b/projects/hip/include/hip/nvcc_detail/hip_runtime.h
@@ -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__
diff --git a/projects/hip/tests/src/kernel/hipLaunchParm.cpp b/projects/hip/tests/src/kernel/hipLaunchParm.cpp
index 992ac5a18f..797bc84ced 100644
--- a/projects/hip/tests/src/kernel/hipLaunchParm.cpp
+++ b/projects/hip/tests/src/kernel/hipLaunchParm.cpp
@@ -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),