From d492f1fd6bab120e22e04d5f2d58b187b547e3dc Mon Sep 17 00:00:00 2001 From: wkwchau Date: Wed, 19 Jun 2019 20:28:29 -0400 Subject: [PATCH] Implement the hipOccupancyMaxPotentialBlockSize function (#1162) * Implement the hipOccupancyMaxPotentialBlockSize function * Replaced hipGetDeviceProperties() call by ihipGetDeviceProperties() in ihipOccupancyMaxPotentialBlockSize() * Add test for hipOccupancyMaxPotentialBlockSize in Module API * Added extern declaration for ihipGetDeviceProperties() to be accessed inside ihipOccupancyMaxPotentialBlockSize() * fixed hipOccupancyMaxPotentialBlockSize test build issue * Fix hipOccupancyMaxPotentialBlockSize dtest * Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest * Revert "Add BUILD_CMD in hipOccupancyMaxPotentialBlockSize dtest" This reverts commit 0480ff56f1441fc515d2c26ce33783e303423938. * Disable hipOccupancyMaxPotentialBlockSize dtest on NVCC * move extern declaration of ihipGetDeviceProperties to hip_module.cpp * Update the limiation of 32 wavefronts per CU and 800/512 SGPRs for VI/pre-VI chips to calculate the occupancy --- .../hip/hcc_detail/functional_grid_launch.hpp | 16 ++ include/hip/hcc_detail/hip_runtime_api.h | 16 +- src/hip_module.cpp | 146 ++++++++++++++++++ .../hipOccupancyMaxPotentialBlockSize.cpp | 69 +++++++++ 4 files changed, 246 insertions(+), 1 deletion(-) create mode 100644 tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index b09d0d2c6c..111a6e2b70 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -135,6 +135,22 @@ void hipLaunchKernelGGLImpl( } } // Namespace hip_impl. + +template +inline +void hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, + F kernel, size_t dynSharedMemPerBlk, uint32_t blockSizeLimit) { + + using namespace hip_impl; + + hip_impl::hip_init(); + auto f = get_program_state().kernel_descriptor(reinterpret_cast(kernel), + target_agent(0)); + + hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f, + dynSharedMemPerBlk, blockSizeLimit); +} + template inline void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks, diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 993fb2b1b4..e70f4954db 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2744,6 +2744,21 @@ hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, dim3 blockDim hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags); +/** + * @brief determine the grid and block sizes to achieves maximum occupancy for a kernel + * + * @param [out] gridSize minimum grid size for maximum potential occupancy + * @param [out] blockSize block size for maximum potential occupancy + * @param [in] f kernel to launch + * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block + * @param [in] blockSizeLimit the maximum block size for the kernel, use 0 for no limit + * + * @returns hipSuccess, hipInvalidDevice, hipErrorInvalidValue + */ +hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, + hipFunction_t f, size_t dynSharedMemPerBlk, + uint32_t blockSizeLimit); + /** * @brief Returns occupancy for a device function. * @@ -2782,7 +2797,6 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, int numDevices, unsigned int flags); - // doxygen end Version Management /** * @} diff --git a/src/hip_module.cpp b/src/hip_module.cpp index d69be7f689..78ecea0488 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -119,6 +119,8 @@ string ToString(hipFunction_t v) { const std::string& FunctionSymbol(const hipFunction_t f) { return f->_name; }; +extern hipError_t ihipGetDeviceProperties(hipDeviceProp_t* props, int device); + #define CHECK_HSA(hsaStatus, hipStatus) \ if (hsaStatus != HSA_STATUS_SUCCESS) { \ return hipStatus; \ @@ -805,3 +807,147 @@ hipError_t hipModuleGetTexRef(textureReference** texRef, hipModule_t hmod, const *texRef = reinterpret_cast(addr); return ihipLogStatus(hipSuccess); } + +hipError_t ihipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, + hipFunction_t f, size_t dynSharedMemPerBlk, + uint32_t blockSizeLimit) +{ + using namespace hip_impl; + + auto ctx = ihipGetTlsDefaultCtx(); + hipError_t ret = hipSuccess; + + if (ctx == nullptr) { + ret = hipErrorInvalidDevice; + } + + hipDeviceProp_t prop{}; + ihipGetDeviceProperties(&prop, ihipGetTlsDefaultCtx()->getDevice()->_deviceId); + + prop.regsPerBlock = prop.regsPerBlock ? prop.regsPerBlock : 64 * 1024; + + size_t usedVGPRS = 0; + size_t usedSGPRS = 0; + size_t usedLDS = 0; + bool is_code_object_v3 = f->_name.find(".kd") != std::string::npos; + if (is_code_object_v3) { + const auto header = reinterpret_cast(f->_header); + // GRANULATED_WAVEFRONT_VGPR_COUNT is specified in 0:5 bits of COMPUTE_PGM_RSRC1 + // the granularity for gfx6-gfx9 is max(0, ceil(vgprs_used / 4) - 1) + usedVGPRS = ((header->compute_pgm_rsrc1 & 0x3F) + 1) << 2; + // GRANULATED_WAVEFRONT_SGPR_COUNT is specified in 6:9 bits of COMPUTE_PGM_RSRC1 + // the granularity for gfx9+ is 2 * max(0, ceil(sgprs_used / 16) - 1) + usedSGPRS = ((((header->compute_pgm_rsrc1 & 0x3C0) >> 6) >> 1) + 1) << 4; + usedLDS = header->group_segment_fixed_size; + } + else { + const auto header = f->_header; + // VGPRs granularity is 4 + usedVGPRS = ((header->workitem_vgpr_count + 3) >> 2) << 2; + // adding 2 to take into account the 2 VCC registers & handle the granularity of 16 + usedSGPRS = header->wavefront_sgpr_count + 2; + usedSGPRS = ((usedSGPRS + 15) >> 4) << 4; + usedLDS = header->workgroup_group_segment_byte_size; + } + + // try different workgroup sizes to find the maximum potential occupancy + // based on the usage of VGPRs and LDS + size_t wavefrontSize = prop.warpSize; + size_t maxWavefrontsPerBlock = prop.maxThreadsPerBlock / wavefrontSize; + + // Due to SPI and private memory limitations, the max of wavefronts per CU in 32 + size_t maxWavefrontsPerCU = min(prop.maxThreadsPerMultiProcessor / wavefrontSize, 32); + + const size_t numSIMD = 4; + size_t maxActivWaves = 0; + size_t maxWavefronts = 0; + for (int i = 0; i < maxWavefrontsPerBlock; i++) { + size_t wavefrontsPerWG = i + 1; + + // workgroup per CU is 40 for WG size of 1 wavefront; otherwise it is 16 + size_t maxWorkgroupPerCU = (wavefrontsPerWG == 1) ? 40 : 16; + size_t maxWavesWGLimited = min(wavefrontsPerWG * maxWorkgroupPerCU, maxWavefrontsPerCU); + + // Compute VGPR limited wavefronts per block + size_t wavefrontsVGPRS; + if (usedVGPRS == 0) { + wavefrontsVGPRS = maxWavesWGLimited; + } + else { + // find how many VGPRs are available for each SIMD + size_t numVGPRsPerSIMD = (prop.regsPerBlock / wavefrontSize / numSIMD); + wavefrontsVGPRS = (numVGPRsPerSIMD / usedVGPRS) * numSIMD; + } + + size_t maxWavesVGPRSLimited = 0; + if (wavefrontsVGPRS > maxWavesWGLimited) { + maxWavesVGPRSLimited = maxWavesWGLimited; + } + else { + maxWavesVGPRSLimited = (wavefrontsVGPRS / wavefrontsPerWG) * wavefrontsPerWG; + } + + // Compute SGPR limited wavefronts per block + size_t wavefrontsSGPRS; + if (usedSGPRS == 0) { + wavefrontsSGPRS = maxWavesWGLimited; + } + else { + const size_t numSGPRsPerSIMD = (prop.gcnArch < 900) ? 512 : 800; + wavefrontsSGPRS = (numSGPRsPerSIMD / usedSGPRS) * numSIMD; + } + + size_t maxWavesSGPRSLimited = 0; + if (wavefrontsSGPRS > maxWavesWGLimited) { + maxWavesSGPRSLimited = maxWavesWGLimited; + } + else { + maxWavesSGPRSLimited = (wavefrontsSGPRS / wavefrontsPerWG) * wavefrontsPerWG; + } + + // Compute LDS limited wavefronts per block + size_t wavefrontsLDS; + if (usedLDS == 0) { + wavefrontsLDS = maxWorkgroupPerCU * wavefrontsPerWG; + } + else { + size_t availableSharedMemPerCU = prop.maxSharedMemoryPerMultiProcessor; + size_t workgroupPerCU = availableSharedMemPerCU / (usedLDS + dynSharedMemPerBlk); + wavefrontsLDS = min(workgroupPerCU, maxWorkgroupPerCU) * wavefrontsPerWG; + } + + size_t maxWavesLDSLimited = min(wavefrontsLDS, maxWavefrontsPerCU); + + size_t activeWavefronts = 0; + size_t tmp_min = (size_t)min(maxWavesLDSLimited, maxWavesWGLimited); + tmp_min = min(maxWavesSGPRSLimited, tmp_min); + activeWavefronts = min(maxWavesVGPRSLimited, tmp_min); + + if (maxActivWaves < activeWavefronts) { + maxActivWaves = activeWavefronts; + maxWavefronts = wavefrontsPerWG; + } + } + + // determine the grid and block sizes for maximum potential occupancy + size_t maxThreadsCnt = prop.maxThreadsPerMultiProcessor*prop.multiProcessorCount; + if (blockSizeLimit > 0) { + maxThreadsCnt = min(maxThreadsCnt, blockSizeLimit); + } + + *blockSize = maxWavefronts * wavefrontSize; + *gridSize = min((maxThreadsCnt + *blockSize - 1) / *blockSize, prop.multiProcessorCount); + + return ret; +} + + +hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* blockSize, + hipFunction_t f, size_t dynSharedMemPerBlk, + uint32_t blockSizeLimit) +{ + HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit); + + return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize( + gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit)); +} diff --git a/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp b/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp new file mode 100644 index 0000000000..a81862952d --- /dev/null +++ b/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp @@ -0,0 +1,69 @@ +/* +Copyright (c) 2019 Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +all copies or substantial portions of the Software. + +THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ +// Test the Grid_Launch syntax. + +/* HIT_START + * BUILD: %t %s ../../test_common.cpp EXCLUDE_HIP_PLATFORM nvcc + * TEST: %t + * HIT_END + */ + +#include "hip/hip_runtime.h" +#include "test_common.h" + +#define fileName "vcpy_kernel.code" +#define kernel_name "hello_world" + + +__global__ void f1(float *a) { *a = 1.0; } + +template +__global__ void f2(T *a) { *a = 1; } + + + +int main(int argc, char* argv[]) { + + // test case for using kernel function pointer + uint32_t gridSize = 0; + uint32_t blockSize = 0; + hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0); + assert(gridSize != 0 && blockSize != 0); + + // test case for using kernel function pointer with template + gridSize = 0; + blockSize = 0; + hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f2, 0, 0); + assert(gridSize != 0 && blockSize != 0); + + // test case for using kernel with hipFunction_t type + gridSize = 0; + blockSize = 0; + hipModule_t Module; + hipFunction_t Function; + HIPCHECK(hipModuleLoad(&Module, fileName)); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + HIPCHECK(hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, Function, 0, 0)); + assert(gridSize != 0 && blockSize != 0); + + passed(); +}