Grid-launch updates to 2.0 and cleanup of old.

_ Use fields from GRID_LAUNCH_20 structure
  (See USE_GRID_LAUNCH_20 define, currently set to 0)
  "1" will require HCC support.
- Remove old DISABLE_GRID_LAUNCH support.

Change-Id: I584ce648d217251789a6283cf27feb24cb7dc8d1

include/hcc_detail/hip_runtime.h

@@ -38,6 +38,10 @@ THE SOFTWARE.
 #include <stddef.h>
 
 
+// Use field names for grid_launch 2.0 structure:
+#define USE_GRID_LAUNCH_20 0
+
+
 #define CUDA_SUCCESS hipSuccess
 
 #include <hip/hip_runtime_api.h>
@@ -517,7 +521,20 @@ void ihipPostLaunchKernel(hipStream_t stream, grid_launch_parm &lp);
 #define KNRM  "\x1B[0m"
 #define KGRN  "\x1B[32m"
 
-#if not defined(DISABLE_GRID_LAUNCH)
+#if USE_GRID_LAUNCH_20
+#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
+do {\
+  grid_launch_parm lp;\
+  lp.dynamic_group_mem_bytes = _groupMemBytes; \
+  hipStream_t trueStream = (ihipPreLaunchKernel(_stream, _numBlocks3D, _blockDim3D, &lp)); \
+    if (HIP_TRACE_API) {\
+        fprintf(stderr, KGRN "<<hip-api: hipLaunchKernel '%s' gridDim:(%d,%d,%d) groupDim:(%d,%d,%d) groupMem:+%d stream=%p\n" KNRM, \
+                #_kernelName, lp.grid_dim.x, lp.grid_dim.y, lp.grid_dim.z, lp.group_dim.x, lp.group_dim.y, lp.group_dim.z, lp.groupMemBytes, (void*)(_stream));\
+    }\
+  _kernelName (lp, ##__VA_ARGS__);\
+  ihipPostLaunchKernel(trueStream, lp);\
+} while(0)
+#else
 #define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
 do {\
   grid_launch_parm lp;\
@@ -528,35 +545,12 @@ do {\
                 #_kernelName, lp.gridDim.x, lp.gridDim.y, lp.gridDim.z, lp.groupDim.x, lp.groupDim.y, lp.groupDim.z, lp.groupMemBytes, (void*)(_stream));\
     }\
   _kernelName (lp, ##__VA_ARGS__);\
-  ihipPostLaunchKernel(trueStream, lp);\
-} while(0)
-
-#else
-#warning(DISABLE_GRID_LAUNCH set)
-
-#define hipLaunchKernel(_kernelName, _numBlocks3D, _blockDim3D, _groupMemBytes, _stream, ...) \
-do {\
-  grid_launch_parm lp;\
-  lp.gridDim.x = _numBlocks3D.x * _blockDim3D.x;/*Convert from #blocks to #threads*/ \
-  lp.gridDim.y = _numBlocks3D.y * _blockDim3D.y;/*Convert from #blocks to #threads*/ \
-  lp.gridDim.z = _numBlocks3D.z * _blockDim3D.z;/*Convert from #blocks to #threads*/ \
-  lp.groupDim.x = _blockDim3D.x; \
-  lp.groupDim.y = _blockDim3D.y; \
-  lp.groupDim.z = _blockDim3D.z; \
-  lp.groupMemBytes = _groupMemBytes;\
-  hc::completion_future cf;\
-  lp.cf = &cf;  \
-  hipStream_t trueStream = (ihipPreLaunchKernel(_stream, &lp.av)); \
-    if (HIP_TRACE_API) {\
-        fprintf(stderr, "==hip-api: launch '%s' gridDim:[%d.%d.%d] groupDim:[%d.%d.%d] groupMem:+%d stream=%p\n", \
-                #_kernelName, lp.gridDim.z, lp.gridDim.y, lp.gridDim.x, lp.groupDim.z, lp.groupDim.y, lp.groupDim.x, lp.groupMemBytes, (void*)(_stream));\
-    }\
-  _kernelName (lp, ##__VA_ARGS__);\
   ihipPostLaunchKernel(trueStream, cf);\
 } while(0)
-/*end hipLaunchKernel */
+
 #endif
 
+
 #elif defined (__HCC_C__)
 
 //TODO - develop C interface.

include/hcc_detail/host_defines.h

@@ -35,11 +35,7 @@ THE SOFTWARE.
 #define __host__     __attribute__((cpu))
 #define __device__   __attribute__((hc))
 
-#ifndef DISABLE_GRID_LAUNCH
 #define __global__  __attribute__((hc_grid_launch))
-#else
-#define __global__
-#endif
 
 #define __noinline__      __attribute__((noinline))
 #define __forceinline__   __attribute__((always_inline))

include/hip_common.h

@@ -22,9 +22,6 @@ THE SOFTWARE.
 
 #pragma once
 
-// Disable use of grid_launch feature in HCC compiler.
-//#define DISABLE_GRID_LAUNCH
-
 // Common code included at start of every hip file.
 // Auto enable __HIP_PLATFORM_HCC__ if compiling with HCC
 // Other compiler (GCC,ICC,etc) need to set one of these macros explicitly