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
This commit is contained in:
wkwchau
Maneesh Gupta
@@ -135,6 +135,22 @@ void hipLaunchKernelGGLImpl(
|
||||
}
|
||||
} // Namespace hip_impl.
|
||||
|
||||
|
||||
template <typename F>
|
||||
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<std::uintptr_t>(kernel),
|
||||
target_agent(0));
|
||||
|
||||
hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f,
|
||||
dynSharedMemPerBlk, blockSizeLimit);
|
||||
}
|
||||
|
||||
template <typename... Args, typename F = void (*)(Args...)>
|
||||
inline
|
||||
void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
|
||||
|
||||
@@ -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
|
||||
/**
|
||||
* @}
|
||||
|
||||
@@ -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<textureReference*>(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<const amd_kernel_code_v3_t*>(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));
|
||||
}
|
||||
|
||||
@@ -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 <typename T>
|
||||
__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<void(*)(int *)>(&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();
|
||||
}
|
||||
Reference in New Issue
Block a user