From 7a3bda3c2f0d58adc24c638c77c13edae05e00ea Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 7 Jun 2019 05:17:15 +0530 Subject: [PATCH 1/8] Fix sample to use kernelargs for launch --- samples/0_Intro/module_api/Makefile | 6 ++--- samples/0_Intro/module_api/defaultDriver.cpp | 28 +++++++++----------- 2 files changed, 16 insertions(+), 18 deletions(-) diff --git a/samples/0_Intro/module_api/Makefile b/samples/0_Intro/module_api/Makefile index 270d4c1211..8d0237826f 100644 --- a/samples/0_Intro/module_api/Makefile +++ b/samples/0_Intro/module_api/Makefile @@ -5,7 +5,7 @@ endif HIPCC=$(HIP_PATH)/bin/hipcc HIP_PLATFORM=$(shell $(HIP_PATH)/bin/hipconfig --compiler) -all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out +all: vcpy_kernel.code runKernel.hip.out launchKernelHcc.hip.out defaultDriver.hip.out runKernel.hip.out: runKernel.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ @@ -13,8 +13,8 @@ runKernel.hip.out: runKernel.cpp launchKernelHcc.hip.out: launchKernelHcc.cpp $(HIPCC) $(HIPCC_FLAGS) $< -o $@ -#defaultDriver.hip.out: defaultDriver.cpp -# $(HIPCC) $(HIPCC_FLAGS) $< -o $@ +defaultDriver.hip.out: defaultDriver.cpp + $(HIPCC) $(HIPCC_FLAGS) $< -o $@ vcpy_kernel.code: vcpy_kernel.cpp $(HIPCC) --genco $(GENCO_FLAGS) $^ -o $@ diff --git a/samples/0_Intro/module_api/defaultDriver.cpp b/samples/0_Intro/module_api/defaultDriver.cpp index b6abfcc3a9..9fa0b0ee09 100644 --- a/samples/0_Intro/module_api/defaultDriver.cpp +++ b/samples/0_Intro/module_api/defaultDriver.cpp @@ -21,7 +21,6 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip/hip_runtime_api.h" #include #include #include @@ -29,20 +28,18 @@ THE SOFTWARE. #define LEN 64 #define SIZE LEN << 2 -#define fileName "test.co" -#define kernel_name "vadd" +#define fileName "vcpy_kernel.code" +#define kernel_name "hello_world" int main() { - float *A, *B, *C; - hipDeviceptr_t Ad, Bd, Cd; + float *A, *B; + hipDeviceptr_t Ad, Bd; A = new float[LEN]; B = new float[LEN]; - C = new float[LEN]; for (uint32_t i = 0; i < LEN; i++) { A[i] = i * 1.0f; - B[i] = 1.0f; - C[i] = 0.0f; + B[i] = 0.0f; } hipInit(0); @@ -53,28 +50,25 @@ int main() { hipMalloc((void**)&Ad, SIZE); hipMalloc((void**)&Bd, SIZE); - hipMalloc((void**)&Cd, SIZE); hipMemcpyHtoD(Ad, A, SIZE); hipMemcpyHtoD(Bd, B, SIZE); - hipMemcpyHtoD(Cd, C, SIZE); hipModule_t Module; hipFunction_t Function; hipModuleLoad(&Module, fileName); hipModuleGetFunction(&Function, Module, kernel_name); - int n = LEN; - void* args[4] = {&Ad, &Bd, &Cd, &n}; + void* args[2] = {&Ad, &Bd}; hipModuleLaunchKernel(Function, 1, 1, 1, LEN, 1, 1, 0, 0, args, nullptr); - hipMemcpyDtoH(C, Cd, SIZE); + hipMemcpyDtoH(B, Bd, SIZE); int mismatchCount = 0; for (uint32_t i = 0; i < LEN; i++) { - if (A[i] + B[i] != C[i]) { + if (A[i] != B[i]) { mismatchCount++; - std::cout << "error: mismatch " << A[i] << " + " << B[i] << " != " << C[i] << std::endl; + std::cout << "error: mismatch " << A[i] << " != " << C[i] << std::endl; } } @@ -84,6 +78,10 @@ int main() { std::cout << "FAILED!\n"; }; + hipFree(Ad); + hipFree(Bd); + delete A; + delete B; hipCtxDestroy(context); return 0; } From f2b3526503fbd7bd81105e838c4efc3418479855 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 7 Jun 2019 05:23:11 +0530 Subject: [PATCH 2/8] Add cleanup code in module api samples --- samples/0_Intro/module_api/launchKernelHcc.cpp | 4 ++++ samples/0_Intro/module_api/runKernel.cpp | 4 ++++ samples/0_Intro/module_api_global/runKernel.cpp | 4 ++++ 3 files changed, 12 insertions(+) diff --git a/samples/0_Intro/module_api/launchKernelHcc.cpp b/samples/0_Intro/module_api/launchKernelHcc.cpp index 5ccd0801d1..4ff000d1af 100644 --- a/samples/0_Intro/module_api/launchKernelHcc.cpp +++ b/samples/0_Intro/module_api/launchKernelHcc.cpp @@ -105,6 +105,10 @@ int main() { std::cout << "FAILED!\n"; }; + hipFree(Ad); + hipFree(Bd); + delete A; + delete B; hipCtxDestroy(context); return 0; } diff --git a/samples/0_Intro/module_api/runKernel.cpp b/samples/0_Intro/module_api/runKernel.cpp index eccc105319..a011b42666 100644 --- a/samples/0_Intro/module_api/runKernel.cpp +++ b/samples/0_Intro/module_api/runKernel.cpp @@ -97,6 +97,10 @@ int main() { std::cout << "FAILED!\n"; }; + hipFree(Ad); + hipFree(Bd); + delete A; + delete B; hipCtxDestroy(context); return 0; } diff --git a/samples/0_Intro/module_api_global/runKernel.cpp b/samples/0_Intro/module_api_global/runKernel.cpp index 76af21b0f9..1e10dba5c6 100644 --- a/samples/0_Intro/module_api_global/runKernel.cpp +++ b/samples/0_Intro/module_api_global/runKernel.cpp @@ -145,6 +145,10 @@ int main() { }; } + hipFree(Ad); + hipFree(Bd); + delete A; + delete B; hipCtxDestroy(context); return 0; } From 19ca7a2a0851b7d852d4fcd763197b9d446dba0a Mon Sep 17 00:00:00 2001 From: Evgeny Date: Tue, 11 Jun 2019 20:13:29 -0500 Subject: [PATCH 3/8] prof layer includes refactoring --- hip_prof_gen.py | 47 ++++++++++--------- include/hip/hcc_detail/hip_runtime_api.h | 6 +-- src/hip_hcc_internal.h | 1 + src/hip_intercept.cpp | 6 ++- .../hip/hcc_detail => src}/hip_prof_api.h | 0 5 files changed, 32 insertions(+), 28 deletions(-) rename {include/hip/hcc_detail => src}/hip_prof_api.h (100%) diff --git a/hip_prof_gen.py b/hip_prof_gen.py index 09bf65ebcd..e17b150cfc 100755 --- a/hip_prof_gen.py +++ b/hip_prof_gen.py @@ -301,8 +301,6 @@ def generate_prof_header(f, api_map, opts_map): f.write('// automatically generated sources\n') f.write('#ifndef _HIP_PROF_STR_H\n'); f.write('#define _HIP_PROF_STR_H\n'); - f.write('#include \n'); - f.write('#include \n'); # Generating dummy macro for non-public API f.write('\n// Dummy API primitives\n') @@ -383,27 +381,30 @@ def generate_prof_header(f, api_map, opts_map): f.write('#define INIT_CB_ARGS_DATA(cb_id, cb_data) INIT_##cb_id##_CB_ARGS_DATA(cb_data)\n') # Generating the method for the API string, name and parameters - f.write('\n') - f.write('#if 0\n') - f.write('// HIP API string method, method name and parameters\n') - f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n') - f.write(' std::ostringstream oss;\n') - f.write(' switch (id) {\n') - for name, args in api_map.items(): - f.write(' case HIP_API_ID_' + name + ':\n') - f.write(' oss << "' + name + '("') - for ind in range(0, len(args)): - arg_tuple = args[ind] - arg_name = arg_tuple[1] - if ind != 0: f.write(' << ","') - f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name) - f.write('\n << ")";\n') - f.write(' break;\n') - f.write(' default: oss << "unknown";\n') - f.write(' };\n') - f.write(' return strdup(oss.str().c_str());\n') - f.write('};\n') - f.write('#endif\n') + if False: + f.write('\n') + f.write('#if 0\n') + f.write('#include \n'); + f.write('#include \n'); + f.write('// HIP API string method, method name and parameters\n') + f.write('const char* hipApiString(hip_api_id_t id, const hip_api_data_t* data) {\n') + f.write(' std::ostringstream oss;\n') + f.write(' switch (id) {\n') + for name, args in api_map.items(): + f.write(' case HIP_API_ID_' + name + ':\n') + f.write(' oss << "' + name + '("') + for ind in range(0, len(args)): + arg_tuple = args[ind] + arg_name = arg_tuple[1] + if ind != 0: f.write(' << ","') + f.write('\n << " ' + arg_name + '=" << data->args.' + name + '.' + arg_name) + f.write('\n << ")";\n') + f.write(' break;\n') + f.write(' default: oss << "unknown";\n') + f.write(' };\n') + f.write(' return strdup(oss.str().c_str());\n') + f.write('};\n') + f.write('#endif\n') f.write('#endif // _HIP_PROF_STR_H\n'); diff --git a/include/hip/hcc_detail/hip_runtime_api.h b/include/hip/hcc_detail/hip_runtime_api.h index 68cefe2410..993fb2b1b4 100644 --- a/include/hip/hcc_detail/hip_runtime_api.h +++ b/include/hip/hcc_detail/hip_runtime_api.h @@ -2973,9 +2973,7 @@ hipError_t hipLaunchByPtr(const void* func); } /* extern "c" */ #endif -#ifdef __cplusplus -#include -#endif +#include #ifdef __cplusplus extern "C" { @@ -2987,7 +2985,7 @@ hipError_t hipRegisterApiCallback(uint32_t id, void* fun, void* arg); hipError_t hipRemoveApiCallback(uint32_t id); hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg); hipError_t hipRemoveActivityCallback(uint32_t id); -static inline const char* hipApiName(const uint32_t& id) { return hip_api_name(id); } +const char* hipApiName(uint32_t id); const char* hipKernelNameRef(const hipFunction_t f); #ifdef __cplusplus } /* extern "C" */ diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index 769d6b7914..7695ea34c1 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -30,6 +30,7 @@ THE SOFTWARE. #include "hsa/hsa_ext_amd.h" #include "hip/hip_runtime.h" +#include "hip_prof_api.h" #include "hip_util.h" #include "env.h" diff --git a/src/hip_intercept.cpp b/src/hip_intercept.cpp index cab8aeb23b..6e8b120360 100644 --- a/src/hip_intercept.cpp +++ b/src/hip_intercept.cpp @@ -21,7 +21,7 @@ THE SOFTWARE. */ #include "hip/hip_runtime.h" -#include "hip/hcc_detail/hip_prof_api.h" +#include "hip_prof_api.h" // HIP API callback/activity @@ -47,3 +47,7 @@ hipError_t hipRegisterActivityCallback(uint32_t id, void* fun, void* arg) { hipError_t hipRemoveActivityCallback(uint32_t id) { return callbacks_table.set_activity(id, NULL, NULL) ? hipSuccess : hipErrorInvalidValue; } + +const char* hipApiName(uint32_t id) { + return hip_api_name(id); +} diff --git a/include/hip/hcc_detail/hip_prof_api.h b/src/hip_prof_api.h similarity index 100% rename from include/hip/hcc_detail/hip_prof_api.h rename to src/hip_prof_api.h From e32940357f3756b5901753569938dce4101558bc Mon Sep 17 00:00:00 2001 From: Michael LIAO Date: Fri, 14 Jun 2019 13:31:51 -0400 Subject: [PATCH 4/8] [hipcc] Revise include path calculation. - Once HIP_VDI_HOME is defined but HIP_CLANG_INCLUDE_PATH is not, calculate it directly without HIP_CLANG_PATH is defined or not; Otherwise, we may leave HIP_CLANG_INCLUDE_PATH undefined, if clang is not installed following the official way (so far, HIP-Clang breaks that), we may leave HIP_CLANG_INCLUDE_PATH undefined before its uses. --- bin/hipcc | 3 +++ 1 file changed, 3 insertions(+) diff --git a/bin/hipcc b/bin/hipcc index 58c8fe45c9..5b25d01fcf 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -118,6 +118,9 @@ if ($HIP_RUNTIME eq "VDI" and !defined $HIP_VDI_HOME) { if (defined $HIP_VDI_HOME) { if (!defined $HIP_CLANG_PATH and (-e "$HIP_VDI_HOME/bin/clang" or -e "$HIP_VDI_HOME/bin/clang.exe")) { $HIP_CLANG_PATH = "$HIP_VDI_HOME/bin"; + } + # With HIP_VDI_HOME defined, assume the installation of clang components, including headers. + if (!defined $HIP_CLANG_INCLUDE_PATH) { $HIP_CLANG_INCLUDE_PATH = "$HIP_VDI_HOME/include/clang"; } if (!defined $DEVICE_LIB_PATH and -e "$HIP_VDI_HOME/lib/bitcode") { From cc374b2bd37ab10fc07edaa2e91f67e71a12a7a3 Mon Sep 17 00:00:00 2001 From: Icarus Sparry <21295823+icarus-sparry@users.noreply.github.com> Date: Mon, 17 Jun 2019 12:03:36 -0700 Subject: [PATCH 5/8] Make hip_prof_gen.py compatible with both python 2 and 3 Convert python 2 constructs to python 3 compatible ones. In python 3, print is a function, so use write methods (which are always functions) instead. In python3 keys() returns an iterator, rather than a list. This means you can't change the data structure that is being iterated over. Converting this iterator into a list mimics the python 2 behavior. --- hip_prof_gen.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/hip_prof_gen.py b/hip_prof_gen.py index 09bf65ebcd..3a2f53ef8e 100755 --- a/hip_prof_gen.py +++ b/hip_prof_gen.py @@ -13,7 +13,7 @@ line_num = -1 # Verbose message def message(msg): - if verbose: print >>sys.stdout, msg + if verbose: sys.stdout.write(msg + '\n') # Fatal error termination def error(msg): @@ -24,8 +24,8 @@ def error(msg): else: msg = " Warning: " + msg - print >>sys.stdout, msg - print >>sys.stderr, sys.argv[0] + msg + sys.stdout.write(msg + '\n') + sys.stderr.write(sys.argv[0] + msg +'\n') def fatal(msg): error(msg) @@ -451,7 +451,7 @@ parse_api(api_hfile, api_map) parse_src(api_map, src_dir, src_pat, opts_map) # Checking for non-conformant APIs -for name in opts_map.keys(): +for name in list(opts_map.keys()): m = re.match(r'\.(\S*)', name) if m: message("Init missing: " + m.group(1)) From d492f1fd6bab120e22e04d5f2d58b187b547e3dc Mon Sep 17 00:00:00 2001 From: wkwchau Date: Wed, 19 Jun 2019 20:28:29 -0400 Subject: [PATCH 6/8] 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(); +} From 96dc74897df03a76559728f54c560efb751ca829 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Wed, 19 Jun 2019 20:29:05 -0400 Subject: [PATCH 7/8] [hip] implement the hipExtLaunchMultiKernelMultiDevice API (#1165) * [hip] implement the hipExtLaunchMultiKernelMultiDevice API * add a guard to check the HCC version for acquire_locked_hsa_queue() API which was introdued in HCC for ROCm 2.5 * modified code based on the requested changes * changes to lock all streams before launching kernels for each device and unlock them after the dispatches * check each stream to be valid before starting to lock all the streams --- include/hip/hcc_detail/hip_runtime.h | 10 +-- include/hip/hcc_detail/program_state.hpp | 1 + src/hip_hcc.cpp | 49 +++++++++------ src/hip_hcc_internal.h | 4 +- src/hip_module.cpp | 78 ++++++++++++++++++++++-- 5 files changed, 113 insertions(+), 29 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 68406f0e29..6ac402a6f7 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -270,14 +270,14 @@ static inline __device__ void printf(const char* format, All... all) {} #if defined __HCC_CPP__ extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, - grid_launch_parm* lp, const char* kernelNameStr); + grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0); extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, - grid_launch_parm* lp, const char* kernelNameStr); + grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0); extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, - grid_launch_parm* lp, const char* kernelNameStr); + grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0); extern hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, - grid_launch_parm* lp, const char* kernelNameStr); -extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp); + grid_launch_parm* lp, const char* kernelNameStr, bool lockAcquired = 0); +extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp, bool unlockPostponed = 0); #if GENERIC_GRID_LAUNCH == 0 //#warning "Original hipLaunchKernel defined" diff --git a/include/hip/hcc_detail/program_state.hpp b/include/hip/hcc_detail/program_state.hpp index c64b64fde8..4e0fbb76a4 100644 --- a/include/hip/hcc_detail/program_state.hpp +++ b/include/hip/hcc_detail/program_state.hpp @@ -77,6 +77,7 @@ class kernargs_size_align { public: std::size_t size(std::size_t n) const; std::size_t alignment(std::size_t n) const; + const void* getHandle() const {return handle;}; private: const void* handle; friend kernargs_size_align program_state::get_kernargs_size_align(std::uintptr_t); diff --git a/src/hip_hcc.cpp b/src/hip_hcc.cpp index 583017a30a..4edf9575be 100644 --- a/src/hip_hcc.cpp +++ b/src/hip_hcc.cpp @@ -404,7 +404,7 @@ LockedAccessor_StreamCrit_t ihipStream_t::lockopen_preKernelCommand() { //--- // Must be called after kernel finishes, this releases the lock on the stream so other commands can // submit. -void ihipStream_t::lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av) { +void ihipStream_t::lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av, bool unlockPostponed) { bool blockThisKernel = false; if (!g_hipLaunchBlockingKernels.empty()) { @@ -426,7 +426,10 @@ void ihipStream_t::lockclose_postKernelCommand(const char* kernelName, hc::accel kernelName); } - _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand. + // if unlockPostponed is true then this stream will be unlocked later (e.g., see hipExtLaunchMultiKernelMultiDevice for a sample call) + if (!unlockPostponed) { + _criticalData.unlock(); // paired with lock from lockopen_preKernelCommand. + } }; @@ -1493,7 +1496,7 @@ void ihipStreamCallbackHandler(ihipStreamCallback_t* cb) { // // If stream==NULL synchronize appropriately with other streams and return the default av for the // device. If stream is valid, return the AV to use. -hipStream_t ihipSyncAndResolveStream(hipStream_t stream) { +hipStream_t ihipSyncAndResolveStream(hipStream_t stream, bool lockAcquired) { if (stream == hipStreamNull) { // Submitting to NULL stream, call locked_syncDefaultStream to wait for all other streams: ihipCtx_t* ctx = ihipGetTlsDefaultCtx(); @@ -1535,9 +1538,14 @@ hipStream_t ihipSyncAndResolveStream(hipStream_t stream) { if (needGatherMarker) { // ensure any commands sent to this stream wait on the NULL stream before // continuing - LockedAccessor_StreamCrit_t thisStreamCrit(stream->criticalData()); - // TODO - could be "noret" version of create_blocking_marker - thisStreamCrit->_av.create_blocking_marker(dcf, hc::accelerator_scope); + if (!lockAcquired) { + LockedAccessor_StreamCrit_t thisStreamCrit(stream->criticalData()); + // TODO - could be "noret" version of create_blocking_marker + thisStreamCrit->_av.create_blocking_marker(dcf, hc::accelerator_scope); + } else { + // this stream is already locked (e.g., call from hipExtLaunchMultiKernelMultiDevice) + stream->criticalData()._av.create_blocking_marker(dcf, hc::accelerator_scope); + } tprintf( DB_SYNC, " %s adding marker to wait for freshly recorded default-stream marker \n", @@ -1578,8 +1586,8 @@ void ihipPrintKernelLaunch(const char* kernelName, const grid_launch_parm* lp, // Called just before a kernel is launched from hipLaunchKernel. // Allows runtime to track some information about the stream. hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_launch_parm* lp, - const char* kernelNameStr) { - stream = ihipSyncAndResolveStream(stream); + const char* kernelNameStr, bool lockAcquired) { + stream = ihipSyncAndResolveStream(stream, lockAcquired); lp->grid_dim.x = grid.x; lp->grid_dim.y = grid.y; lp->grid_dim.z = grid.z; @@ -1589,8 +1597,13 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ lp->barrier_bit = barrier_bit_queue_default; lp->launch_fence = -1; - auto crit = stream->lockopen_preKernelCommand(); - lp->av = &(crit->_av); + if (!lockAcquired) { + auto crit = stream->lockopen_preKernelCommand(); + lp->av = &(crit->_av); + } else { + // this stream is already locked (e.g., call from hipExtLaunchMultiKernelMultiDevice) + lp->av = &(stream->criticalData()._av); + } lp->cf = nullptr; ihipPrintKernelLaunch(kernelNameStr, lp, stream); @@ -1599,30 +1612,30 @@ hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, dim3 block, grid_ hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, dim3 block, grid_launch_parm* lp, - const char* kernelNameStr) { - return ihipPreLaunchKernel(stream, dim3(grid), block, lp, kernelNameStr); + const char* kernelNameStr, bool lockAcquired) { + return ihipPreLaunchKernel(stream, dim3(grid), block, lp, kernelNameStr, lockAcquired); } hipStream_t ihipPreLaunchKernel(hipStream_t stream, dim3 grid, size_t block, grid_launch_parm* lp, - const char* kernelNameStr) { - return ihipPreLaunchKernel(stream, grid, dim3(block), lp, kernelNameStr); + const char* kernelNameStr, bool lockAcquired) { + return ihipPreLaunchKernel(stream, grid, dim3(block), lp, kernelNameStr, lockAcquired); } hipStream_t ihipPreLaunchKernel(hipStream_t stream, size_t grid, size_t block, grid_launch_parm* lp, - const char* kernelNameStr) { - return ihipPreLaunchKernel(stream, dim3(grid), dim3(block), lp, kernelNameStr); + const char* kernelNameStr, bool lockAcquired) { + return ihipPreLaunchKernel(stream, dim3(grid), dim3(block), lp, kernelNameStr, lockAcquired); } //--- // Called after kernel finishes execution. // This releases the lock on the stream. -void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp) { +void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, grid_launch_parm& lp, bool unlockPostponed) { tprintf(DB_SYNC, "ihipPostLaunchKernel, unlocking stream\n"); - stream->lockclose_postKernelCommand(kernelName, lp.av); + stream->lockclose_postKernelCommand(kernelName, lp.av, unlockPostponed); if (HIP_PROFILE_API) { MARKER_END(); } diff --git a/src/hip_hcc_internal.h b/src/hip_hcc_internal.h index 7695ea34c1..852e8985ad 100644 --- a/src/hip_hcc_internal.h +++ b/src/hip_hcc_internal.h @@ -550,7 +550,7 @@ class ihipStream_t { // Member functions that begin with locked_ are thread-safe accessors - these acquire / release // the critical mutex. LockedAccessor_StreamCrit_t lockopen_preKernelCommand(); - void lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av); + void lockclose_postKernelCommand(const char* kernelName, hc::accelerator_view* av, bool unlockNotNeeded = 0); void locked_wait(); @@ -952,7 +952,7 @@ hipError_t hipModuleGetFunctionEx(hipFunction_t* hfunc, hipModule_t hmod, const char* name, hsa_agent_t *agent); -hipStream_t ihipSyncAndResolveStream(hipStream_t); +hipStream_t ihipSyncAndResolveStream(hipStream_t, bool lockAcquired = 0); hipError_t ihipStreamSynchronize(hipStream_t stream); void ihipStreamCallbackHandler(ihipStreamCallback_t* cb); diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 78ecea0488..bda6970298 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -150,7 +150,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, uint32_t localWorkSizeX, uint32_t localWorkSizeY, uint32_t localWorkSizeZ, size_t sharedMemBytes, hipStream_t hStream, void** kernelParams, void** extra, - hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags) { + hipEvent_t startEvent, hipEvent_t stopEvent, uint32_t flags, bool isStreamLocked = 0) { using namespace hip_impl; auto ctx = ihipGetTlsDefaultCtx(); @@ -206,8 +206,7 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, sharedMemBytes; // TODO - this should be part of preLaunchKernel. hStream = ihipPreLaunchKernel( hStream, dim3(globalWorkSizeX/localWorkSizeX, globalWorkSizeY/localWorkSizeY, globalWorkSizeZ/localWorkSizeZ), - dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str()); - + dim3(localWorkSizeX, localWorkSizeY, localWorkSizeZ), &lp, f->_name.c_str(), isStreamLocked); hsa_kernel_dispatch_packet_t aql; @@ -272,7 +271,9 @@ hipError_t ihipModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, stopEvent->attachToCompletionFuture(&cf, hStream, hipEventTypeStopCommand); } - ihipPostLaunchKernel(f->_name.c_str(), hStream, lp); + ihipPostLaunchKernel(f->_name.c_str(), hStream, lp, isStreamLocked); + + } return ret; @@ -315,6 +316,75 @@ hipError_t hipHccModuleLaunchKernel(hipFunction_t f, uint32_t globalWorkSizeX, localWorkSizeZ, sharedMemBytes, hStream, kernelParams, extra, startEvent, stopEvent, 0)); } +hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, + int numDevices, unsigned int flags) { + + hipError_t result; + + if ((numDevices > g_deviceCnt) || (launchParamsList == nullptr)) { + return hipErrorInvalidValue; + } + + hipFunction_t* kds = reinterpret_cast(malloc(sizeof(hipFunction_t) * numDevices)); + if (kds == nullptr) { + return hipErrorNotInitialized; + } + + // prepare all kernel descriptors for each device as all streams will be locked in the next loop + for (int i = 0; i < numDevices; ++i) { + const hipLaunchParams& lp = launchParamsList[i]; + if (lp.stream == nullptr) { + free(kds); + return hipErrorNotInitialized; + } + kds[i] = hip_impl::get_program_state().kernel_descriptor(reinterpret_cast(lp.func), + hip_impl::target_agent(lp.stream)); + if (kds[i] == nullptr) { + free(kds); + return hipErrorInvalidValue; + } + hip_impl::kernargs_size_align kargs = hip_impl::get_program_state().get_kernargs_size_align( + reinterpret_cast(lp.func)); + kds[i]->_kernarg_layout = *reinterpret_cast>*>( + kargs.getHandle()); + } + + // lock all streams before launching kernels to each device + for (int i = 0; i < numDevices; ++i) { + LockedAccessor_StreamCrit_t streamCrit(launchParamsList[i].stream->criticalData(), false); + #if (__hcc_workweek__ >= 19213) + streamCrit->_av.acquire_locked_hsa_queue(); + #endif + } + + // launch kernels for each device + for (int i = 0; i < numDevices; ++i) { + const hipLaunchParams& lp = launchParamsList[i]; + + result = ihipModuleLaunchKernel(kds[i], + lp.gridDim.x * lp.blockDim.x, + lp.gridDim.y * lp.blockDim.y, + lp.gridDim.z * lp.blockDim.z, + lp.blockDim.x, lp.blockDim.y, + lp.blockDim.z, lp.sharedMem, + lp.stream, lp.args, nullptr, nullptr, nullptr, 0, + true /* stream is already locked above and will be unlocked + in the below code after launching kernels on all devices*/); + } + + // unlock all streams + for (int i = 0; i < numDevices; ++i) { + launchParamsList[i].stream->criticalData().unlock(); + #if (__hcc_workweek__ >= 19213) + launchParamsList[i].stream->criticalData()._av.release_locked_hsa_queue(); + #endif + } + + free(kds); + + return result; +} + namespace hip_impl { hsa_executable_t executable_for(hipModule_t hmod) { return hmod->executable; From db6571ae750c3a38c44740e4714b4798c0c49956 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 20 Jun 2019 18:05:57 +0300 Subject: [PATCH 8/8] [HIPIFY] Fix multiple input files support Reported in #1168 --- hipify-clang/src/main.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 25b21bf56f..368d21a3d4 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -234,7 +234,7 @@ int main(int argc, const char **argv) { argc++; } llcompat::PrintStackTraceOnErrorSignal(); - ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::Optional); + ct::CommonOptionsParser OptionsParser(argc, argv, ToolTemplateCategory, llvm::cl::ZeroOrMore); std::vector fileSources = OptionsParser.getSourcePathList(); if (fileSources.empty() && !GeneratePerl && !GeneratePython) { llvm::errs() << "\n" << sHipify << sError << "Must specify at least 1 positional argument for source file." << "\n";