diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index da57240ebd..c493eec933 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -127,6 +127,35 @@ void hipLaunchKernelGGLImpl( } // Namespace hip_impl. +template +inline +hipError_t 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)); + + return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f, + dynSharedMemPerBlk, blockSizeLimit); +} + +template +inline +hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(uint32_t* numBlocks, F kernel, + uint32_t blockSize, size_t dynSharedMemPerBlk) { + + using namespace hip_impl; + + hip_impl::hip_init(); + auto f = get_program_state().kernel_descriptor(reinterpret_cast(kernel), + target_agent(0)); + + return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f, blockSize, dynSharedMemPerBlk); +} + 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 624f3615e1..f2518c3b75 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2923,14 +2923,14 @@ hipError_t hipLaunchCooperativeKernelMultiDevice(hipLaunchParams* launchParamsLi * @param [out] gridSize minimum grid size for maximum potential occupancy * @param [out] blockSize block size for maximum potential occupancy * @param [in] f kernel function for which occupancy is calulated - * @param [in] dynamicSMemSize Per - block dynamic shared memory usage intended, in bytes + * @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(int* gridSize, int* blockSize, - const void* f, size_t dynamicSMemSize, - int blockSizeLimit); +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. @@ -2938,10 +2938,10 @@ hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, * @param [out] numBlocks Returned occupancy * @param [in] func Kernel function for which occupancy is calulated * @param [in] blockSize Block size the kernel is intended to be launched with - * @param [in] dynamicSMemSize Per - block dynamic shared memory usage intended, in bytes + * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block */ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( - int* numBlocks, const void* f, int blockSize, size_t dynamicSMemSize); + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk); /** * @brief Returns occupancy for a device function. @@ -2949,11 +2949,11 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( * @param [out] numBlocks Returned occupancy * @param [in] func Kernel function for which occupancy is calulated * @param [in] blockSize Block size the kernel is intended to be launched with - * @param [in] dynamicSMemSize Per - block dynamic shared memory usage intended, in bytes + * @param [in] dynSharedMemPerBlk dynamic shared memory usage (in bytes) intended for each block * @param [in] flags Extra flags for occupancy calculation (currently ignored) */ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, const void* f, int blockSize, size_t dynamicSMemSize, unsigned int flags); + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags); /** * @brief Launches kernels on multiple devices and guarantees all specified kernels are dispatched @@ -3359,27 +3359,7 @@ hipError_t hipBindTextureToMipmappedArray(const texture& tex, return hipSuccess; } -template -inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( - int* numBlocks, T f, int blockSize, size_t dynamicSMemSize) { - return hipOccupancyMaxActiveBlocksPerMultiprocessor( - numBlocks, reinterpret_cast(f), blockSize, dynamicSMemSize); -} -template -inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, T f, int blockSize, size_t dynamicSMemSize, unsigned int flags) { - return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - numBlocks, reinterpret_cast(f), blockSize, dynamicSMemSize, flags); -} - -template -inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, - T f, size_t dynamicSMemSize, int blockSizeLimit) { - return hipOccupancyMaxPotentialBlockSize( - gridSize, blockSize, reinterpret_cast(f), dynamicSMemSize, blockSizeLimit); -} - template inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim, void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) { diff --git a/samples/2_Cookbook/13_occupancy/occupancy.cpp b/samples/2_Cookbook/13_occupancy/occupancy.cpp index 01fa7aafed..a9f4e198b0 100644 --- a/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -56,9 +56,9 @@ void launchKernel(float* C, float* A, float* B, bool manual){ const unsigned threadsperblock = 32; const unsigned blocks = (NUM/threadsperblock)+1; - int mingridSize = 0; - int gridSize = 0; - int blockSize = 0; + uint32_t mingridSize = 0; + uint32_t gridSize = 0; + uint32_t blockSize = 0; if (manual){ blockSize = threadsperblock; @@ -86,7 +86,7 @@ void launchKernel(float* C, float* A, float* B, bool manual){ printf("kernel Execution time = %6.3fms\n", eventMs); //Calculate Occupancy - int numBlock = 0; + uint32_t numBlock = 0; HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0)); if(devProp.maxThreadsPerMultiProcessor){ diff --git a/src/hip_module.cpp b/src/hip_module.cpp index c15cc34cd5..2afbabf0a8 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -1248,9 +1248,9 @@ void getGprsLdsUsage(hipFunction_t f, size_t* usedVGPRS, size_t* usedSGPRS, size } } -hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, int* gridSize, int* blockSize, - hipFunction_t f, size_t dynamicSMemSize, - int blockSizeLimit) +hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, uint32_t* gridSize, uint32_t* blockSize, + hipFunction_t f, size_t dynSharedMemPerBlk, + uint32_t blockSizeLimit) { using namespace hip_impl; @@ -1331,7 +1331,7 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, int* gridSize, int* } else { size_t availableSharedMemPerCU = prop.maxSharedMemoryPerMultiProcessor; - size_t workgroupPerCU = availableSharedMemPerCU / (usedLDS + dynamicSMemSize); + size_t workgroupPerCU = availableSharedMemPerCU / (usedLDS + dynSharedMemPerBlk); wavefrontsLDS = min(workgroupPerCU, maxWorkgroupPerCU) * wavefrontsPerWG; } @@ -1360,19 +1360,18 @@ hipError_t ihipOccupancyMaxPotentialBlockSize(TlsData *tls, int* gridSize, int* return hipSuccess; } -hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize, - const void* f, size_t dynamicSMemSize, - int blockSizeLimit) +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, dynamicSMemSize, blockSizeLimit); - auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f), - hip_impl::target_agent(0)); + HIP_INIT_API(hipOccupancyMaxPotentialBlockSize, gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit); + return ihipLogStatus(ihipOccupancyMaxPotentialBlockSize(tls, - gridSize, blockSize, F, dynamicSMemSize, blockSizeLimit)); + gridSize, blockSize, f, dynSharedMemPerBlk, blockSizeLimit)); } hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( - TlsData *tls, int* numBlocks, hipFunction_t f, int blockSize, size_t dynamicSMemSize) + TlsData *tls, uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk) { using namespace hip_impl; @@ -1412,39 +1411,35 @@ hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( : std::min(maxWavesPerSimd, availableSGPRs / usedSGPRS)); // Calculate blocks occupancy per CU based on SGPR usage - *numBlocks = std::min(*numBlocks, (int) (sgprs_alu_occupancy / numWavefronts)); + *numBlocks = std::min(*numBlocks, (uint32_t) (sgprs_alu_occupancy / numWavefronts)); - size_t total_used_lds = usedLDS + dynamicSMemSize; + size_t total_used_lds = usedLDS + dynSharedMemPerBlk; if (total_used_lds != 0) { // Calculate LDS occupacy per CU. lds_per_cu / (static_lsd + dynamic_lds) size_t lds_occupancy = prop.maxSharedMemoryPerMultiProcessor / total_used_lds; - *numBlocks = std::min(*numBlocks, (int) lds_occupancy); + *numBlocks = std::min(*numBlocks, (uint32_t) lds_occupancy); } return hipSuccess; } hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( - int* numBlocks, const void* f, int blockSize, size_t dynamicSMemSize) + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk) { - HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynamicSMemSize); + HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessor, numBlocks, f, blockSize, dynSharedMemPerBlk); - auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f), - hip_impl::target_agent(0)); return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( - tls, numBlocks, F, blockSize, dynamicSMemSize)); + tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, const void* f, int blockSize, size_t dynamicSMemSize, + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags) { - HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynamicSMemSize, flags); + HIP_INIT_API(hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags, numBlocks, f, blockSize, dynSharedMemPerBlk, flags); - auto F = hip_impl::get_program_state().kernel_descriptor((std::uintptr_t)(f), - hip_impl::target_agent(0)); return ihipLogStatus(ihipOccupancyMaxActiveBlocksPerMultiprocessor( - tls, numBlocks, F, blockSize, dynamicSMemSize)); + tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } hipError_t hipLaunchKernel( diff --git a/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp b/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp index 8089f26f16..89d003ea94 100644 --- a/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp +++ b/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp @@ -116,7 +116,7 @@ int main() { dimBlock.x = workgroups[i]; // Calculate the device occupancy to know how many blocks can be run concurrently - hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlocks, + hipOccupancyMaxActiveBlocksPerMultiprocessor(reinterpret_cast(&numBlocks), test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long)); dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32); diff --git a/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp b/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp index 2838c09cd1..ebf656b72f 100644 --- a/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp +++ b/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp @@ -30,6 +30,10 @@ THE SOFTWARE. #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 @@ -40,15 +44,16 @@ __global__ void f2(T *a) { *a = 1; } int main(int argc, char* argv[]) { // test case for using kernel function pointer - int gridSize = 0; - int blockSize = 0; + uint32_t gridSize = 0; + uint32_t blockSize = 0; hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0); assert(gridSize != 0 && blockSize != 0); - int numBlock = 0; + uint32_t numBlock = 0; hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, blockSize, 0); assert(numBlock != 0); + // test case for using kernel function pointer with template gridSize = 0; blockSize = 0; @@ -59,5 +64,15 @@ int main(int argc, char* argv[]) { hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f2, blockSize, 0); assert(numBlock != 0); + + // test case for using kernel with hipFunction_t type + numBlock = 0; + hipModule_t Module; + hipFunction_t Function; + HIPCHECK(hipModuleLoad(&Module, fileName)); + HIPCHECK(hipModuleGetFunction(&Function, Module, kernel_name)); + HIPCHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, Function, blockSize, 0)); + assert(numBlock != 0); + passed(); } diff --git a/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp b/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp index 22a3f05283..a81862952d 100644 --- a/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp +++ b/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp @@ -30,16 +30,22 @@ THE SOFTWARE. #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 - int gridSize = 0; - int blockSize = 0; + uint32_t gridSize = 0; + uint32_t blockSize = 0; hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0); assert(gridSize != 0 && blockSize != 0); @@ -49,5 +55,15 @@ int main(int argc, char* argv[]) { 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(); }