From e23cd0dd3c76f422a3869241c758abd59dd80863 Mon Sep 17 00:00:00 2001
From: Ben Sander <ben.sander@amd.com>
Date: Tue, 9 Aug 2016 15:36:50 -0500
Subject: [PATCH] Document workaround for parenthesis+macro+hipLaunchKernel

Change-Id: Ie04c99db92d6499ddde93028a96f9d8f72d3f992
---
 hipamd/docs/markdown/hip_porting_guide.md | 32 +++++++++++++
 hipamd/tests/src/hipLaunchParm.cpp        | 55 +++++++++++++++++++++++
 2 files changed, 87 insertions(+)

diff --git a/hipamd/docs/markdown/hip_porting_guide.md b/hipamd/docs/markdown/hip_porting_guide.md
index 7857e4b983..4d376f4a5f 100644
--- a/hipamd/docs/markdown/hip_porting_guide.md
+++ b/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
diff --git a/hipamd/tests/src/hipLaunchParm.cpp b/hipamd/tests/src/hipLaunchParm.cpp
index c6d28fcd3a..2f4bf11ea2 100644
--- a/hipamd/tests/src/hipLaunchParm.cpp
+++ b/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();
 }