From e27b5cc927d9de41d35aa247f058dfb95148be18 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Mon, 20 Jun 2016 23:46:42 -0500 Subject: [PATCH] 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 | 31 --------------- hipamd/include/hcc_detail/hip_runtime.h | 46 ++++++++++------------- hipamd/include/hcc_detail/host_defines.h | 4 -- hipamd/include/hip_common.h | 3 -- hipamd/src/hip_hcc.cpp | 44 ++++++++++++++++++++++ hipamd/tests/src/kernel/hipGridLaunch.cpp | 1 - 6 files changed, 64 insertions(+), 65 deletions(-) diff --git a/hipamd/docs/markdown/hip_porting_guide.md b/hipamd/docs/markdown/hip_porting_guide.md index de5c590e12..6ae0a6d31e 100644 --- a/hipamd/docs/markdown/hip_porting_guide.md +++ b/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), diff --git a/hipamd/include/hcc_detail/hip_runtime.h b/hipamd/include/hcc_detail/hip_runtime.h index 066db4bbec..8594954e75 100644 --- a/hipamd/include/hcc_detail/hip_runtime.h +++ b/hipamd/include/hcc_detail/hip_runtime.h @@ -38,6 +38,10 @@ THE SOFTWARE. #include +// Use field names for grid_launch 2.0 structure: +#define USE_GRID_LAUNCH_20 0 + + #define CUDA_SUCCESS hipSuccess #include @@ -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 "<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; diff --git a/hipamd/tests/src/kernel/hipGridLaunch.cpp b/hipamd/tests/src/kernel/hipGridLaunch.cpp index f13781362e..b195a0171d 100644 --- a/hipamd/tests/src/kernel/hipGridLaunch.cpp +++ b/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"