Document workaround for parenthesis+macro+hipLaunchKernel

Change-Id: Ie04c99db92d6499ddde93028a96f9d8f72d3f992

hipamd/docs/markdown/hip_porting_guide.md

@@ -264,6 +264,38 @@ Makefiles can use the following syntax to conditionally provide a default HIP_PA
 HIP_PATH ?= $(shell hipconfig --path)
 ```
 
+## hipLaunchKernel 
+
+hipLaunchKernel is a variadic macro which accepts as parameters the launch configurations (grid dims, group dims, stream, dynamic shared size) followed by a variable number of kernel arguments.
+This sequence is then expanded into the appropriate kernel launch syntax depending on the platform.  
+While this can be a convenient single-line kernel launch syntax, the macro implementation can cause issues when nested inside other macros.  For example, consider the following:
+
+```
+// Will cause compile error:
+#define MY_LAUNCH(command, doTrace) \
+{\
+    if (doTrace) printf ("TRACE: %s\n", #command); \
+    (command);   /* The nested ( ) will cause compile error */\
+}
+
+MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall");
+```
+
+Avoid nesting macro parameters inside parenthesis - here's an alternative that will work:
+
+```
+#define MY_LAUNCH(command, doTrace) \
+{\
+    if (doTrace) printf ("TRACE: %s\n", #command); \
+    command;\ 
+}
+
+MY_LAUNCH (hipLaunchKernel(vAdd, dim3(1024), dim3(1), 0, 0, Ad), true, "firstCall");
+```
+
+
+
+
 ## Compiler Options
 
 hipcc is a portable compiler driver that will call nvcc or hcc (depending on the target system) and attach all required include and library options. It passes options through to the target compiler. Tools that call hipcc must ensure the compiler options are appropriate for the target compiler. The `hipconfig` script may helpful in making

hipamd/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();
 }