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

hipamd/docs/markdown/hip_porting_guide.md

@@ -464,37 +464,6 @@ hipcc-cmd: /opt/hcc/bin/hcc  -hc -I/opt/hcc/include -stdlib=libc++ -I../../../..
 If you pass a ".cu" file, hcc will attempt to compile it as a Cuda language file. You must tell hcc that it’s in fact a C++ file: use the "-x c++" option.
 
 
-#### grid_launch kernel dispatch - fallback
-HIP uses an hcc language feature called "grid_launch". The [[hc_grid_launch]] attribute that can be attached to a function definition, and the first parameter is of type grid_launch_parm.
-When a [[hc_grid_launch]] function is called, hcc runtime uses the grid_launch_parm to control the execution configuration of the kernel 
-(including the grid and group dimensions, the queue, and dynamic group memory allocations).   By default, the hipLaunchKernel macro creates a grid_launch_parm structure and launches a
-[[hc_grid_launch]] kernel.  grid_launch is a relatively new addition to hcc so this section describes how to fall back to a traditional calling sequence which invokes a standard host function
-which calls a hc::parallel_for_each to launch the kernel.  
-
-First, set DISABLE_GRID_LAUNCH:
-include/hip_common.h
-```
-// Set this define to disable GRID_LAUNCH
-#define DISABLE_GRID_LAUNCH
-```
-
-Inside any kernel use the KERNELBEGIN as the first line in the kernel function, and KERNELEND as the last line.  For example:
-```
-__global__ void
-MyKernel(hipLaunchParm lp, float *C, const float *A, size_t N)
-{
-    KERNELBEGIN; // Required if hc_grid_launch is disabled
-
-	int tid = hipBlockIdx_x*MAX_THREADS_PER_BLOCK + hipThreadIdx_x;
-
-    if (tid < N) {
-        C[tid] = A[tid];
-    }
-
-    KERNELEND; // Required if hc_grid_launch is disabled
-}
-```
-
 #### HIP Environment Variables
 
 On the HCC path, HIP provides a number of environment variables that control the behavior of HIP.  Some of these are useful for appliction development (for example HIP_VISIBLE_DEVICES, HIP_LAUNCH_BLOCKING),

hipamd/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.

hipamd/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))

hipamd/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

hipamd/src/hip_hcc.cpp

@@ -1091,12 +1091,23 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_
 {
 	std::call_once(hip_initialized, ihipInit);
     stream = ihipSyncAndResolveStream(stream);
+#if USE_GRID_LAUNCH_20 
+    lp->grid_dim.x = grid.x;
+    lp->grid_dim.y = grid.y;
+    lp->grid_dim.z = grid.z;
+    lp->group_dim.x = block.x;
+    lp->group_dim.y = block.y;
+    lp->group_dim.z = block.z;
+    lp->barrier_bit = barrier_bit_queue_default;
+    lp->launch_fence = -1;
+#else
     lp->gridDim.x = grid.x;
     lp->gridDim.y = grid.y;
     lp->gridDim.z = grid.z;
     lp->groupDim.x = block.x;
     lp->groupDim.y = block.y;
     lp->groupDim.z = block.z;
+#endif
     stream->lockopen_preKernelCommand();
 //    *av = &stream->_av;
     lp->av = &stream->_av;
@@ -1109,12 +1120,23 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, gri
 {
 	std::call_once(hip_initialized, ihipInit);
     stream = ihipSyncAndResolveStream(stream);
+#if USE_GRID_LAUNCH_20 
+    lp->grid_dim.x = grid;
+    lp->grid_dim.y = 1;
+    lp->grid_dim.z = 1;
+    lp->group_dim.x = block.x;
+    lp->group_dim.y = block.y;
+    lp->group_dim.z = block.z;
+    lp->barrier_bit = barrier_bit_queue_default;
+    lp->launch_fence = -1;
+#else
     lp->gridDim.x = grid;
     lp->gridDim.y = 1;
     lp->gridDim.z = 1;
     lp->groupDim.x = block.x;
     lp->groupDim.y = block.y;
     lp->groupDim.z = block.z;
+#endif
     stream->lockopen_preKernelCommand();
 //    *av = &stream->_av;
     lp->av = &stream->_av;
@@ -1128,12 +1150,23 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, gri
 {
 	std::call_once(hip_initialized, ihipInit);
     stream = ihipSyncAndResolveStream(stream);
+#if USE_GRID_LAUNCH_20 
+    lp->grid_dim.x = grid.x;
+    lp->grid_dim.y = grid.y;
+    lp->grid_dim.z = grid.z;
+    lp->group_dim.x = block;
+    lp->group_dim.y = 1;
+    lp->group_dim.z = 1;
+    lp->barrier_bit = barrier_bit_queue_default;
+    lp->launch_fence = -1;
+#else
     lp->gridDim.x = grid.x;
     lp->gridDim.y = grid.y;
     lp->gridDim.z = grid.z;
     lp->groupDim.x = block;
     lp->groupDim.y = 1;
     lp->groupDim.z = 1;
+#endif
     stream->lockopen_preKernelCommand();
 //    *av = &stream->_av;
     lp->av = &stream->_av;
@@ -1147,12 +1180,23 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, g
 {
 	std::call_once(hip_initialized, ihipInit);
     stream = ihipSyncAndResolveStream(stream);
+#if USE_GRID_LAUNCH_20 
+    lp->grid_dim.x = grid;
+    lp->grid_dim.y = 1;
+    lp->grid_dim.z = 1;
+    lp->group_dim.x = block;
+    lp->group_dim.y = 1;
+    lp->group_dim.z = 1;
+    lp->barrier_bit = barrier_bit_queue_default;
+    lp->launch_fence = -1;
+#else
     lp->gridDim.x = grid;
     lp->gridDim.y = 1;
     lp->gridDim.z = 1;
     lp->groupDim.x = block;
     lp->groupDim.y = 1;
     lp->groupDim.z = 1;
+#endif
     stream->lockopen_preKernelCommand();
 //    *av = &stream->_av;
     lp->av = &stream->_av;

hipamd/tests/src/kernel/hipGridLaunch.cpp

@@ -21,7 +21,6 @@ THE SOFTWARE.
 */
 // Test the Grid_Launch syntax.
 
-#undef DISABLE_GRID_LAUNCH /* Tell hip_*.h to compile in GL mode */
 #include "hip_runtime.h"
 #include "test_common.h"