diff --git a/include/hip/hcc_detail/functional_grid_launch.hpp b/include/hip/hcc_detail/functional_grid_launch.hpp index 9eb738cf04..76a04fa355 100644 --- a/include/hip/hcc_detail/functional_grid_launch.hpp +++ b/include/hip/hcc_detail/functional_grid_launch.hpp @@ -154,6 +154,20 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block 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 659a6c3c3a..0712db17f9 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2948,7 +2948,7 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block * @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 dynSharedMemPerBlk); + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk); /** * @brief Returns occupancy for a device function. @@ -2960,7 +2960,7 @@ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( * @param [in] flags Extra flags for occupancy calculation (currently ignored) */ hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags); + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags); #if __HIP_VDI__ && !defined(__HCC__) /** @@ -3230,6 +3230,21 @@ hipError_t hipLaunchKernel(const void* function_address, } /* extern "c" */ #endif +#if defined(__cplusplus) && !defined(__HCC__) && defined(__clang__) && defined(__HIP__) +template +static hipError_t __host__ inline hipOccupancyMaxActiveBlocksPerMultiprocessor( + uint32_t* numBlocks, F func, uint32_t blockSize, size_t dynSharedMemPerBlk) { + return ::hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, (hipFunction_t)func, blockSize, + dynSharedMemPerBlk); +} +template +static hipError_t __host__ inline hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + uint32_t* numBlocks, F func, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int flags) { + return ::hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( + numBlocks, (hipFunction_t)func, blockSize, dynSharedMemPerBlk, flags); +} +#endif // defined(__cplusplus) && !defined(__HCC__) && defined(__clang__) && defined(__HIP__) + #if USE_PROF_API #include #endif @@ -3370,20 +3385,6 @@ hipError_t hipBindTextureToMipmappedArray(const texture& tex, return hipSuccess; } -template -inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor( - int* numBlocks, T f, int blockSize, size_t dynSharedMemPerBlk) { - return hipOccupancyMaxActiveBlocksPerMultiprocessor( - numBlocks, reinterpret_cast(f), blockSize, dynSharedMemPerBlk); -} - -template -inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, T f, int blockSize, size_t dynSharedMemPerBlk, unsigned int flags) { - return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - numBlocks, reinterpret_cast(f), blockSize, dynSharedMemPerBlk, flags); -} - #if __HIP_VDI__ && !defined(__HCC__) template inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim, diff --git a/samples/2_Cookbook/13_occupancy/occupancy.cpp b/samples/2_Cookbook/13_occupancy/occupancy.cpp index 605c7724b2..a9f4e198b0 100644 --- a/samples/2_Cookbook/13_occupancy/occupancy.cpp +++ b/samples/2_Cookbook/13_occupancy/occupancy.cpp @@ -86,8 +86,8 @@ void launchKernel(float* C, float* A, float* B, bool manual){ printf("kernel Execution time = %6.3fms\n", eventMs); //Calculate Occupancy - int numBlock = 0; - HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply,(int)blockSize, 0)); + uint32_t numBlock = 0; + HIP_CHECK(hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, multiply, blockSize, 0)); if(devProp.maxThreadsPerMultiProcessor){ std::cout << "Theoretical Occupancy is " << (double)numBlock* blockSize/devProp.maxThreadsPerMultiProcessor * 100 << "%" << std::endl; diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 7d4c97a7fd..16606e8016 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -1385,7 +1385,7 @@ hipError_t hipOccupancyMaxPotentialBlockSize(uint32_t* gridSize, uint32_t* block } hipError_t ihipOccupancyMaxActiveBlocksPerMultiprocessor( - TlsData *tls, int* numBlocks, hipFunction_t f, int blockSize, size_t dynSharedMemPerBlk) + TlsData *tls, uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk) { using namespace hip_impl; @@ -1425,41 +1425,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 + 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 dynSharedMemPerBlk) + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk) { 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, dynSharedMemPerBlk)); + tls, numBlocks, f, blockSize, dynSharedMemPerBlk)); } hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags( - int* numBlocks, const void* f, int blockSize, size_t dynSharedMemPerBlk, + uint32_t* numBlocks, hipFunction_t f, uint32_t blockSize, size_t dynSharedMemPerBlk, unsigned int 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, dynSharedMemPerBlk)); + 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 896738892d..c76685fa89 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/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp b/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp similarity index 80% rename from tests/src/runtimeApi/occupancy/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp rename to tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp index 8e0dd033bc..ebf656b72f 100644 --- a/tests/src/runtimeApi/occupancy/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 @@ -45,10 +49,11 @@ int main(int argc, char* argv[]) { hipOccupancyMaxPotentialBlockSize(&gridSize, &blockSize, f1, 0, 0); assert(gridSize != 0 && blockSize != 0); - int numBlock = 0; - hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f1, (int)blockSize, 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; @@ -56,7 +61,17 @@ int main(int argc, char* argv[]) { assert(gridSize != 0 && blockSize != 0); numBlock = 0; - hipOccupancyMaxActiveBlocksPerMultiprocessor(&numBlock, f2, (int)blockSize, 0); + 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/occupancy/hipOccupancyMaxPotentialBlockSize.cpp b/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp similarity index 99% rename from tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp rename to tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp index d29100d9a9..a81862952d 100644 --- a/tests/src/runtimeApi/occupancy/hipOccupancyMaxPotentialBlockSize.cpp +++ b/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp @@ -33,6 +33,7 @@ THE SOFTWARE. #define fileName "vcpy_kernel.code" #define kernel_name "hello_world" + __global__ void f1(float *a) { *a = 1.0; } template