Greatly enhance handling of macros in kernel launches

All but the most contrived use of macros is now properly handled -
have a look at the new testcases this commit adds. You can have
macros in kernel calls, macros spanning chunks of your arguments,
the call, call parameters, or callee can all be macros or
partially macros.


[ROCm/hip commit: 094b2b9b05]

projects/hip/tests/hipify-clang/axpy.cu

@@ -2,11 +2,23 @@
 
 #include <iostream>
 
-__global__ void axpy(float a, float* x, float* y) {
+
+#define TOKEN_PASTE(X, Y) X ## Y
+#define ARG_LIST_AS_MACRO a, device_x, device_y
+#define KERNEL_CALL_AS_MACRO axpy<float><<<1, kDataLen>>>
+#define KERNEL_NAME_MACRO axpy<float>
+
+// CHECK: #define COMPLETE_LAUNCH hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y)
+#define COMPLETE_LAUNCH axpy<<<1, kDataLen>>>(a, device_x, device_y)
+
+
+template<typename T>
+__global__ void axpy(T a, T *x, T *y) {
   // CHECK: y[hipThreadIdx_x] = a * x[hipThreadIdx_x];
   y[threadIdx.x] = a * x[threadIdx.x];
 }
 
+
 int main(int argc, char* argv[]) {
   const int kDataLen = 4;
 
@@ -27,10 +39,29 @@ int main(int argc, char* argv[]) {
   // CHECK: hipMemcpy(device_x, host_x, kDataLen * sizeof(float), hipMemcpyHostToDevice);
   cudaMemcpy(device_x, host_x, kDataLen * sizeof(float), cudaMemcpyHostToDevice);
 
-  // Launch the kernel.
+  // Launch the kernel in numerous different strange ways to exercise the prerocessor.
   // CHECK: hipLaunchKernelGGL(axpy, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
   axpy<<<1, kDataLen>>>(a, device_x, device_y);
 
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, device_x, device_y);
+  axpy<float><<<1, kDataLen>>>(a, device_x, device_y);
+
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, a, TOKEN_PASTE(device, _x), device_y);
+  axpy<float><<<1, kDataLen>>>(a, TOKEN_PASTE(device, _x), device_y);
+
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
+  axpy<float><<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
+
+  // CHECK: hipLaunchKernelGGL(KERNEL_NAME_MACRO, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
+  KERNEL_NAME_MACRO<<<1, kDataLen>>>(ARG_LIST_AS_MACRO);
+
+  // CHECK: hipLaunchKernelGGL(axpy<float>, dim3(1), dim3(kDataLen), 0, 0, ARG_LIST_AS_MACRO);
+  KERNEL_CALL_AS_MACRO(ARG_LIST_AS_MACRO);
+
+  // CHECK: COMPLETE_LAUNCH;
+  COMPLETE_LAUNCH;
+
+
   // Copy output data to host.
   // CHECK: hipDeviceSynchronize();
   cudaDeviceSynchronize();