Document workaround for parenthesis+macro+hipLaunchKernel

Change-Id: Ie04c99db92d6499ddde93028a96f9d8f72d3f992

tests/src/hipLaunchParm.cpp

@@ -24,6 +24,39 @@ THE SOFTWARE.
 
 __global__ void vAdd(hipLaunchParm lp, float *a){}
 
+
+//---
+//Some wrapper macro for testing:
+#define WRAP(...) __VA_ARGS__
+
+#include <sys/time.h>
+#define GPU_PRINT_TIME(cmd, elapsed, quiet) do {\
+        struct timeval start, stop;\
+        float elapsed;\
+        gettimeofday(&start, NULL);\
+        hipDeviceSynchronize();\
+        cmd;\
+        hipDeviceSynchronize();\
+        gettimeofday(&stop, NULL);\
+    } while(0); 
+
+
+
+#define MY_LAUNCH(command, doTrace, msg) \
+{\
+    if (doTrace) printf ("TRACE: %s %s\n", msg, #command); \
+    command;\
+}
+
+
+#define MY_LAUNCH_WITH_PAREN(command, doTrace, msg) \
+{\
+    if (doTrace) printf ("TRACE: %s %s\n", msg, #command); \
+    (command);\
+}
+
+
+
 int main()
 {
     float *Ad;
@@ -32,5 +65,27 @@ int main()
     hipLaunchKernel(vAdd, 1024, dim3(1), 0, 0, Ad);
     hipLaunchKernel(vAdd, dim3(1024), 1, 0, 0, Ad);
     hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad);
+
+    // Test case with hipLaunchKernel inside another macro:
+    float e0;
+    GPU_PRINT_TIME (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), e0, j);
+    GPU_PRINT_TIME (WRAP(hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j);
+
+#ifdef EXTRA_PARENS_1
+    // Don't wrap hipLaunchKernel in extra set of parens:
+    GPU_PRINT_TIME ((hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad)), e0, j);
+#endif
+
+    MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall");
+
+    float *A;
+    float e1;
+    MY_LAUNCH_WITH_PAREN (hipMalloc(&A, 100), true, "launch2");
+
+#ifdef EXTRA_PARENS_2
+    //MY_LAUNCH_WITH_PAREN wraps cmd in () which can cause issues.
+    MY_LAUNCH_WITH_PAREN (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall");
+#endif
+
     passed();
 }