AI/rocm-systems
2
0
Fork 0
Código Issues Pull Requests Ações Pacotes Projetos Versões Wiki Atividade

Revert "Fix occupany APIs (#1560)"

Ver Código-Fonte 
This reverts commit 6c5fbf9b4a.
Esse commit está contido em:
Rahul Garg
2019-10-29 11:41:08 -07:00
commit de GitHub
pai 8a7e6fb747
commit 27221bc823
7 arquivos alterados com 98 adições e 63 exclusões
Baixar Arquivo de Patch Baixar Arquivo de Diff Expandir todos os arquivos Colapsar todos os arquivos
hipamd/include/hip/hcc_detail/functional_grid_launch.hpp
+29
Ver Arquivo
@@ -127,6 +127,35 @@ void hipLaunchKernelGGLImpl(
} // Namespace hip_impl.
template <typename F>
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<std::uintptr_t>(kernel),
target_agent(0));
return hipOccupancyMaxPotentialBlockSize(gridSize, blockSize, f,
dynSharedMemPerBlk, blockSizeLimit);
}
template <typename F>
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<std::uintptr_t>(kernel),
target_agent(0));
return hipOccupancyMaxActiveBlocksPerMultiprocessor(numBlocks, f, blockSize, dynSharedMemPerBlk);
}
template <typename... Args, typename F = void (*)(Args...)>
inline
void hipLaunchKernelGGL(F kernel, const dim3& numBlocks, const dim3& dimBlocks,
hipamd/include/hip/hcc_detail/hip_runtime_api.h
+8 -28
Ver Arquivo
@@ -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<T, dim, readMode>& tex,
return hipSuccess;
}
template <class T>
inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessor(
int* numBlocks, T f, int blockSize, size_t dynamicSMemSize) {
return hipOccupancyMaxActiveBlocksPerMultiprocessor(
numBlocks, reinterpret_cast<const void*>(f), blockSize, dynamicSMemSize);
}
template <class T>
inline hipError_t hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
int* numBlocks, T f, int blockSize, size_t dynamicSMemSize, unsigned int flags) {
return hipOccupancyMaxActiveBlocksPerMultiprocessorWithFlags(
numBlocks, reinterpret_cast<const void*>(f), blockSize, dynamicSMemSize, flags);
}
template <class T>
inline hipError_t hipOccupancyMaxPotentialBlockSize(int* gridSize, int* blockSize,
T f, size_t dynamicSMemSize, int blockSizeLimit) {
return hipOccupancyMaxPotentialBlockSize(
gridSize, blockSize, reinterpret_cast<const void*>(f), dynamicSMemSize, blockSizeLimit);
}
template <class T>
inline hipError_t hipLaunchCooperativeKernel(T f, dim3 gridDim, dim3 blockDim,
void** kernelParams, unsigned int sharedMemBytes, hipStream_t stream) {
hipamd/samples/2_Cookbook/13_occupancy/occupancy.cpp
+4 -4
Ver Arquivo
@@ -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){
hipamd/src/hip_module.cpp
+20 -25
Ver Arquivo
@@ -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(
hipamd/tests/src/runtimeApi/module/hipLaunchCooperativeKernel.cpp
+1 -1
Ver Arquivo
@@ -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<uint32_t*>(&numBlocks),
test_gws, dimBlock.x * dimBlock.y * dimBlock.z, dimBlock.x * sizeof(long));
dimGrid.x = deviceProp.multiProcessorCount * std::min(numBlocks, 32);
hipamd/tests/src/runtimeApi/module/hipOccupancyMaxActiveBlocksPerMultiprocessor.cpp
+18 -3
Ver Arquivo
@@ -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 <typename T>
@@ -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<void(*)(int *)>(&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();
}
hipamd/tests/src/runtimeApi/module/hipOccupancyMaxPotentialBlockSize.cpp
+18 -2
Ver Arquivo
@@ -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 <typename T>
__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<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();
}