From a9488603b43d47b478a9a9af901ecfdf52ef2717 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 15 Dec 2016 21:00:34 +0300 Subject: [PATCH 01/12] [HIPIFY] nested macro is not hipified, when it isAnyIdentifier Fix for https://github.com/GPUOpen-ProfessionalCompute-Tools/HIP/issues/55 --- hipamd/hipify-clang/src/Cuda2Hip.cpp | 17 +++++++++++++---- 1 file changed, 13 insertions(+), 4 deletions(-) diff --git a/hipamd/hipify-clang/src/Cuda2Hip.cpp b/hipamd/hipify-clang/src/Cuda2Hip.cpp index 4f37bb9365..064f4ad4e1 100644 --- a/hipamd/hipify-clang/src/Cuda2Hip.cpp +++ b/hipamd/hipify-clang/src/Cuda2Hip.cpp @@ -1465,11 +1465,11 @@ public: // to workaround the 'const' MacroArgs passed into this hook. const Token *start = Args->getUnexpArgument(i); size_t len = Args->getArgLength(start) + 1; - #if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) +#if (LLVM_VERSION_MAJOR >= 3) && (LLVM_VERSION_MINOR >= 9) _pp->EnterTokenStream(ArrayRef(start, len), false); - #else +#else _pp->EnterTokenStream(start, len, false, false); - #endif +#endif do { toks.push_back(Token()); Token &tk = toks.back(); @@ -1489,8 +1489,17 @@ public: << " found as an actual argument in expansion of macro " << macroName << "\n" << "will be replaced with: " << repName << "\n"); + size_t length = name.size(); SourceLocation sl = tok.getLocation(); - Replacement Rep(*_sm, sl, name.size(), repName); + if (_sm->isMacroBodyExpansion(sl)) { + LangOptions DefaultLangOptions; + SourceLocation sl_macro = _sm->getExpansionLoc(sl); + SourceLocation sl_end = Lexer::getLocForEndOfToken(sl_macro, 0, *_sm, DefaultLangOptions); + length = _sm->getCharacterData(sl_end) - _sm->getCharacterData(sl_macro); + name = StringRef(_sm->getCharacterData(sl_macro), length); + sl = sl_macro; + } + Replacement Rep(*_sm, sl, length, repName); Replace->insert(Rep); } } else if (tok.isLiteral()) { From a30b36afdc5d645d9ad95bb1b772cde8e249b100 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 15 Dec 2016 14:41:27 -0600 Subject: [PATCH 02/12] remove TODO file --- hipamd/samples/1_Utils/hipCommander/TODO | 50 ------------------------ 1 file changed, 50 deletions(-) delete mode 100644 hipamd/samples/1_Utils/hipCommander/TODO diff --git a/hipamd/samples/1_Utils/hipCommander/TODO b/hipamd/samples/1_Utils/hipCommander/TODO deleted file mode 100644 index 4c835cfced..0000000000 --- a/hipamd/samples/1_Utils/hipCommander/TODO +++ /dev/null @@ -1,50 +0,0 @@ -_ Add AQL kernel. -_ Fix &*kernel command so the kernel name/type is an argument not a new command. - -_ Add command to parse only. -_ Add regression to parse all the hcm files. - -_ Partition HCC, HIP, HSA, OpenCL commands into separate files. - - -_ Show time for back-to-back copies. -_ Add variables. - %loopcnt - - ./hipCommander %loopcnt=4 - -_ Add datasize command. - - -_ Add ( ) to parsing. -_ Add argument parsing and checking. - -_ Add verbose option to print each step of setup. - - print deliniater between setup and run. Add run start message. - - - print sizes of all buffers. - - print each command before running. - - show start/stop of timer routine. - -_ -_ Clear documentation on what each oepration does. -_ Add time instrumentation for each command. -_ Add pcie atomic. - - -_ Add tests for negative cases, ie endloop w/o opening loop. - - -README tips ---- -- HIP_API_TRACE combined with -v is useful to track the exact commands generates by hipCommander. - - -Other ideas: ---- -[ ] Perf guide : stream creation very slow on HCC and should be avoided. - - -Scratch: - - From b53ba931e8be0cc545e835ff7885b043ae21c0d1 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 15 Dec 2016 14:42:02 -0600 Subject: [PATCH 03/12] fix copyright --- hipamd/src/device_util.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/hipamd/src/device_util.cpp b/hipamd/src/device_util.cpp index 7efb12d2d0..5452e1c905 100644 --- a/hipamd/src/device_util.cpp +++ b/hipamd/src/device_util.cpp @@ -14,7 +14,8 @@ 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, +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. */ From d39ef58dafc3540696a646bfc1c504e7b99e5325 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Thu, 15 Dec 2016 14:42:27 -0600 Subject: [PATCH 04/12] Fix typo --- hipamd/docs/markdown/hip_profiling.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/hipamd/docs/markdown/hip_profiling.md b/hipamd/docs/markdown/hip_profiling.md index 21133100ec..0c55acf85e 100644 --- a/hipamd/docs/markdown/hip_profiling.md +++ b/hipamd/docs/markdown/hip_profiling.md @@ -335,7 +335,7 @@ HIP provides 3 environment variables in the HIP_*_BLOCKING family. These introd These options cause HCC to serialize. Useful if you have libraries or code which is calling HCC kernels directly rather than using HIP. - HCC_SERIALZIE_KERNELS : 0x1=pre-serialize before each kernel launch, 0x2=post-serialize after each kernel launch., 0x3= pre- and post- serialize. -- HCC_SERIALIZE_COPY : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize.0 +- HCC_SERIALIZE_COPY : 0x1=pre-serialize before each async copy, 0x2=post-serialize after each async copy., 0x3= pre- and post- serialize. - HSA_ENABLE_SDMA=0 : Causes host-to-device and device-to-host copies to use compute shader blit kernels rather than the dedicated DMA copy engines. Compute shader copies have low latency (typically < 5us) and can achieve approximately 80% of the bandwidth of the DMA copy engine. This flag is useful to isolate issues with the hardware copy engines. - HSA_ENABLE_INTERRUPT=0 : Causes completion signals to be detected with memory-based polling rather than interrupts. Can be useful to diagnose interrupt storm issues in the driver. From ab07d9b4a77c8be3779b481e3feb949833a23869 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Fri, 16 Dec 2016 08:55:11 -0600 Subject: [PATCH 05/12] Print limits on CUDA devices --- hipamd/samples/1_Utils/hipInfo/hipInfo.cpp | 19 +++++++++++++++++++ 1 file changed, 19 insertions(+) diff --git a/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp b/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp index 42a879e732..cf4660eae7 100644 --- a/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/hipamd/samples/1_Utils/hipInfo/hipInfo.cpp @@ -63,6 +63,14 @@ double bytesToGB(size_t s) return (double)s / (1024.0*1024.0*1024.0); } +#define printLimit(w1, limit, units) \ +{\ + size_t val;\ + cudaDeviceGetLimit(&val, limit);\ + std::cout << setw(w1) << #limit": " << val << " " << units << std::endl;\ +} + + void printDeviceProp (int deviceId) { using namespace std; @@ -144,6 +152,17 @@ void printDeviceProp (int deviceId) cout << endl; +#ifdef __HIP_PLATFORM_NVCC__ + // Limits: + cout << endl; + printLimit(w1, cudaLimitStackSize, "bytes/thread"); + printLimit(w1, cudaLimitPrintfFifoSize, "bytes/device"); + printLimit(w1, cudaLimitMallocHeapSize, "bytes/device"); + printLimit(w1, cudaLimitDevRuntimeSyncDepth, "grids"); + printLimit(w1, cudaLimitDevRuntimePendingLaunchCount, "launches"); +#endif + + cout << endl; From c673aec9715277b8a8cec44513cfaa3e5f75df9c Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Fri, 16 Dec 2016 09:24:59 -0600 Subject: [PATCH 06/12] disabled half native support as inline asm is not working Change-Id: I3073d8ae39eed321987f0f2f0e689eec4cdbb48c --- hipamd/include/hip/hcc_detail/hip_fp16.h | 4 +-- hipamd/src/hip_fp16.cpp | 3 -- hipamd/src/hip_ir.ll | 45 ------------------------ 3 files changed, 1 insertion(+), 51 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h index fb6cd0a44c..d51a5d1fcd 100644 --- a/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -25,9 +25,7 @@ THE SOFTWARE. #include "hip/hip_runtime.h" -#define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__ - -#ifdef HIP_HALF_HW_SUPPORT +#if 0 typedef __fp16 __half; diff --git a/hipamd/src/hip_fp16.cpp b/hipamd/src/hip_fp16.cpp index 3bf6bd395f..1a9d04474f 100644 --- a/hipamd/src/hip_fp16.cpp +++ b/hipamd/src/hip_fp16.cpp @@ -22,8 +22,6 @@ THE SOFTWARE. #include"hip/hip_fp16.h" -#if __CLANG_VERSION__ == 35 - static const unsigned sign_val = 0x8000; static const __half __half_value_one_float = {0x3C00}; static const __half __half_value_zero_float = {0x0}; @@ -374,4 +372,3 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){ __device__ __half2 __low2half2(const __half2 a, const __half2 b){ return {a.q, b.q}; } -#endif diff --git a/hipamd/src/hip_ir.ll b/hipamd/src/hip_ir.ll index 4b7bd3b10e..472038df6a 100644 --- a/hipamd/src/hip_ir.ll +++ b/hipamd/src/hip_ir.ll @@ -12,51 +12,6 @@ define linkonce_odr spir_func void @__threadfence_block() #1 { ret void } -define linkonce_odr spir_func i32 @__rocm_dp4a(i32 %in1, i32 %in2, i32 %in3) { - %val1 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_0 src1_sel:BYTE_0","=v,v,v"(i32 %in1, i32 %in2) - %ret1 = add i32 %val1, %in3 - %val2 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_1 src1_sel:BYTE_1","=v,v,v"(i32 %in1, i32 %in2) - %ret2 = add i32 %ret1, %val2 - %val3 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_2 src1_sel:BYTE_2","=v,v,v"(i32 %in1, i32 %in2) - %ret3 = add i32 %val3, %ret2 - %val4 = tail call i32 asm "v_mul_u32_u24_sdwa $0, $1, $2 dst_sel:DWORD dst_unused:UNUSED_PAD src0_sel:BYTE_3 src1_sel:BYTE_3","=v,v,v"(i32 %in1, i32 %in2) - %ret4 = add i32 %val4, %ret3 - ret i32 %ret4 -} -define linkonce_odr spir_func i32 @__rocm_hfma(i32 %in1, i32 %in2, i32 %in3) { - tail call void asm "v_mac_f16 $0, $1, $2","v,v,v"(i32 %in1, i32 %in2, i32 %in3) - ret i32 %in3 -} - -define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) { - %val = tail call i32 asm "v_add_f16 $0, $1, $2","=v,v,v"(i32 %in1, i32 %in2) - ret i32 %val -} - -define linkonce_odr spir_func half @__hip_hadd_gfx803(half %a, half %b) #1 { - %val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b) - ret half %val -} - -define linkonce_odr spir_func half @__hip_hfma_gfx803(half %a, half %b, half %c) #1 { - %val = tail call half asm "v_fma_f16 $0, $1, $2, $3","=v,v,v,v"(half %a, half %b, half %c) - ret half %val -} - -define linkonce_odr spir_func half @__hip_hmul_gfx803(half %a, half %b) #1 { - %val = tail call half asm "v_mul_f16 $0, $1, $2","=v,v,v"(half %a, half %b) - ret half %val -} - -define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 { - %val = tail call half asm "v_sub_f16 $0, $1, $2","=v,v,v"(half %a, half %b) - ret half %val -} - -define linkonce_odr spir_func i32 @__hip_hadd2_gfx803(i32 %a, i32 %b) #1 { - %val = tail call i32 asm "v_add_f16_sdwa $0, $1, $2 dst_sel:WORD_0 dst_unused:UNUSED_PRESERVE src0_sel:WORD_0 src1_sel:WORD_0","=v,v,v"(i32 %a, i32 %b) - ret i32 %val -} attributes #1 = { alwaysinline nounwind } From bddaa0e81cf2e6095fd157b5caf06ce57511862d Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Sat, 17 Dec 2016 16:53:03 +0530 Subject: [PATCH 07/12] Mapped hipDevice_t to int Change-Id: I6cfa56c42b7cd04aa0e0bce510c0d72d34ea211a --- .../include/hip/hcc_detail/hip_runtime_api.h | 14 +-------- hipamd/src/hip_context.cpp | 14 +++++---- hipamd/src/hip_device.cpp | 30 +++++-------------- hipamd/tests/src/hipDrvGetPCIBusId.cpp | 3 +- 4 files changed, 20 insertions(+), 41 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 52e59ed17a..b2cfa59d71 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -53,7 +53,7 @@ extern "C" { typedef struct ihipCtx_t *hipCtx_t; // Note many APIs also use integer deviceIds as an alternative to the device pointer: -typedef struct ihipDevice_t *hipDevice_t; +typedef int hipDevice_t; typedef struct ihipStream_t *hipStream_t; @@ -1904,18 +1904,6 @@ hipError_t hipIpcCloseMemHandle(void *devPtr); } /* extern "c" */ #endif -#ifdef __cplusplus -/** - * @brief Returns a PCI Bus Id string for the device. - * @param [out] pciBusId - * @param [in] len - * @param [hipDevice_t] device - * - * @returns #hipSuccess, #hipErrorInavlidDevice - */ -hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device); -#endif - /** *------------------------------------------------------------------------------------------------- *------------------------------------------------------------------------------------------------- diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp index 8a2451671c..b38a6c3b74 100644 --- a/hipamd/src/hip_context.cpp +++ b/hipamd/src/hip_context.cpp @@ -57,8 +57,8 @@ hipError_t hipCtxCreate(hipCtx_t *ctx, unsigned int flags, hipDevice_t device) { HIP_INIT_API(ctx, flags, device); // FIXME - review if we want to init hipError_t e = hipSuccess; - - *ctx = new ihipCtx_t(device, g_deviceCnt, flags); + auto deviceHandle = ihipGetDevice(device); + *ctx = new ihipCtx_t(deviceHandle, g_deviceCnt, flags); ihipSetTlsDefaultCtx(*ctx); tls_ctxStack.push(*ctx); @@ -69,11 +69,13 @@ hipError_t hipDeviceGet(hipDevice_t *device, int deviceId) { HIP_INIT_API(device, deviceId); // FIXME - review if we want to init - *device = ihipGetDevice(deviceId); + auto deviceHandle = ihipGetDevice(deviceId); hipError_t e = hipSuccess; - if (*device == NULL) { + if (deviceHandle == NULL) { e = hipErrorInvalidDevice; + } else { + *device = deviceId; } return ihipLogStatus(e); @@ -199,9 +201,11 @@ hipError_t hipCtxGetDevice(hipDevice_t *device) if(ctx == nullptr) { e = hipErrorInvalidContext; + // TODO *device = nullptr; } else { - *device = (ihipDevice_t*)ctx->getDevice(); + auto deviceHandle = ctx->getDevice(); + *device = deviceHandle->_deviceId; } return ihipLogStatus(e); } diff --git a/hipamd/src/hip_device.cpp b/hipamd/src/hip_device.cpp index bf45125b60..1cfdaa619d 100644 --- a/hipamd/src/hip_device.cpp +++ b/hipamd/src/hip_device.cpp @@ -321,9 +321,8 @@ hipError_t hipDeviceComputeCapability(int *major, int *minor, hipDevice_t device { HIP_INIT_API(major,minor, device); hipError_t e = hipSuccess; - int deviceId= device->_deviceId; - e = ihipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, deviceId); - e = ihipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, deviceId); + e = ihipDeviceGetAttribute(major, hipDeviceAttributeComputeCapabilityMajor, device); + e = ihipDeviceGetAttribute(minor, hipDeviceAttributeComputeCapabilityMinor, device); return ihipLogStatus(e); } @@ -331,28 +330,13 @@ hipError_t hipDeviceGetName(char *name,int len,hipDevice_t device) { HIP_INIT_API(name,len, device); hipError_t e = hipSuccess; - int nameLen = strlen(device->_props.name); + auto deviceHandle = ihipGetDevice(device); + int nameLen = strlen(deviceHandle->_props.name); if(nameLen <= len) - memcpy(name,device->_props.name,nameLen); + memcpy(name,deviceHandle->_props.name,nameLen); return ihipLogStatus(e); } -#ifdef __cplusplus -hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len,hipDevice_t device) -{ - HIP_INIT_API(pciBusId, len, device); - hipError_t e = hipSuccess; - int deviceId= device->_deviceId; - int tempPciBusId = 0; - e = ihipDeviceGetAttribute( &tempPciBusId, hipDeviceAttributePciBusId, deviceId); - if( e == hipSuccess) { - std::string tempPciStr = std::to_string(tempPciBusId); - memcpy( pciBusId , tempPciStr.c_str() , tempPciStr.length() ); - } - return ihipLogStatus(e); -} -#endif - hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device) { HIP_INIT_API(pciBusId, len, device); @@ -365,11 +349,13 @@ hipError_t hipDeviceGetPCIBusId (char *pciBusId,int len, int device) } return ihipLogStatus(e); } + hipError_t hipDeviceTotalMem (size_t *bytes,hipDevice_t device) { HIP_INIT_API(bytes, device); hipError_t e = hipSuccess; - *bytes= device->_props.totalGlobalMem; + auto deviceHandle = ihipGetDevice(device); + *bytes= deviceHandle->_props.totalGlobalMem; return ihipLogStatus(e); } diff --git a/hipamd/tests/src/hipDrvGetPCIBusId.cpp b/hipamd/tests/src/hipDrvGetPCIBusId.cpp index 21c49c194c..15c86b2214 100644 --- a/hipamd/tests/src/hipDrvGetPCIBusId.cpp +++ b/hipamd/tests/src/hipDrvGetPCIBusId.cpp @@ -27,7 +27,8 @@ int main(){ hipInit(0); hipDevice_t device; hipDeviceGet(&device, 0); - char pciBusId[100]; + char pciBusId[10]; + memset(pciBusId,0,10); hipDeviceGetPCIBusId(pciBusId,100,device); printf("PCI Bus ID= %s\n",pciBusId); return 0; From 3f9404d0e13c4fd5d137f3478533019a6bd0ca06 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 17 Dec 2016 07:19:22 -0600 Subject: [PATCH 08/12] Refactor Module and Function APIs. - hipFunction_t is now returned by value. This eliminates dynamic allocation / memory management complexity in the module. Removed the kernel name so the structure is just 16 bytes now. - Moved the hsa_executable_load_module and hsa_executable_freeze calls to the hipModuleLoad and hipModuleLoadData calls. - Apply sharedMemBytes in hipModuleLaunchKernel to group segment size (not private). --- .../include/hip/hcc_detail/hip_runtime_api.h | 9 ++- hipamd/src/hip_hcc.h | 35 +--------- hipamd/src/hip_module.cpp | 64 ++++++++++--------- hipamd/src/trace_helper.h | 9 +++ 4 files changed, 52 insertions(+), 65 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index b2cfa59d71..e38f1d6982 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -33,7 +33,7 @@ THE SOFTWARE. #include #include -//#include "hip/hip_hcc.h" +#include #if defined (__HCC__) && (__hcc_workweek__ < 16155) #error("This version of HIP requires a newer version of HCC."); @@ -72,7 +72,12 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t; typedef struct ihipModule_t *hipModule_t; -typedef struct ihipFunction_t *hipFunction_t; +struct ihipModuleSymbol_t{ + hsa_executable_symbol_t _hsaSymbol; + uint64_t _object; // The kernel object. +}; + +typedef struct ihipModuleSymbol_t hipFunction_t; typedef void* hipDeviceptr_t; diff --git a/hipamd/src/hip_hcc.h b/hipamd/src/hip_hcc.h index b5417080e0..4075c2b068 100644 --- a/hipamd/src/hip_hcc.h +++ b/hipamd/src/hip_hcc.h @@ -383,26 +383,6 @@ public: size_t psize; }; -class ihipFunction_t{ -public: - ihipFunction_t(const char *name) { - size_t nameSz = strlen(name); - char *kernelName = (char*)malloc(nameSz); - strncpy(kernelName, name, nameSz); - _kernelName = kernelName; - }; - - ~ihipFunction_t() { - if (_kernelName) { - free((void*)_kernelName); - _kernelName = NULL; - }; - }; -public: - const char *_kernelName; - hsa_executable_symbol_t _kernelSymbol; - uint64_t _kernel; -}; class ihipModule_t { public: @@ -412,20 +392,7 @@ public: void *ptr; size_t size; - ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0), hipFunctionTable() {} - ~ihipModule_t() { - for (int i = 0; i < hipFunctionTable.size(); ++i) { - ihipFunction_t *func = hipFunctionTable[i]; - delete func; - } - hipFunctionTable.clear(); - } - - void registerFunction(ihipFunction_t* func) { - hipFunctionTable.push_back(func); - } -private: - std::vector hipFunctionTable; + ihipModule_t() : executable(), object(), fileName(), ptr(nullptr), size(0) {} }; template diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index badeac2dcd..cb56fb3b4a 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -35,6 +35,11 @@ THE SOFTWARE. //TODO Use Pool APIs from HCC to get memory regions. +#define CHECK_HSA(hsaStatus, hipStatus) \ +if (hsaStatus != HSA_STATUS_SUCCESS) {\ + return ihipLogStatus(hipStatus);\ +} + namespace hipdrv { hsa_status_t findSystemRegions(hsa_region_t region, void *data){ @@ -154,9 +159,13 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_freeze((*module)->executable, NULL); + CHECK_HSA(status, hipErrorNotInitialized); } } @@ -189,7 +198,7 @@ hipError_t hipModuleUnload(hipModule_t hmod) return ihipLogStatus(ret); } -hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const char *name){ +hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char *name){ auto ctx = ihipGetTlsDefaultCtx(); hipError_t ret = hipSuccess; @@ -201,27 +210,21 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch ret = hipErrorInvalidContext; }else{ - *func = new ihipFunction_t(name); - hmod->registerFunction(*func); int deviceId = ctx->getDevice()->_deviceId; ihipDevice_t *currentDevice = ihipGetDevice(deviceId); hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; - hsa_status_t status; - status = hsa_executable_load_code_object(hmod->executable, gpuAgent, hmod->object, NULL); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } - status = hsa_executable_freeze(hmod->executable, NULL); - status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(*func)->_kernelSymbol); + + hsa_status_t status; + status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(func->_hsaSymbol)); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); } - status = hsa_executable_symbol_get_info((*func)->_kernelSymbol, + status = hsa_executable_symbol_get_info(func->_hsaSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, - &(*func)->_kernel); + &func->_object); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); @@ -234,7 +237,7 @@ hipError_t ihipModuleGetFunction(hipFunction_t *func, hipModule_t hmod, const ch hipError_t hipModuleGetFunction(hipFunction_t *hfunc, hipModule_t hmod, const char *name){ HIP_INIT_API(hfunc, hmod, name); - return ihipModuleGetFunction(hfunc, hmod, name); + return ihipModuleGetSymbol(hfunc, hmod, name); } @@ -275,7 +278,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, } uint32_t groupSegmentSize; - hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol, + hsa_status_t status = hsa_executable_symbol_get_info(f._hsaSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, &groupSegmentSize); if(status != HSA_STATUS_SUCCESS){ @@ -283,20 +286,19 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, } uint32_t privateSegmentSize; - status = hsa_executable_symbol_get_info(f->_kernelSymbol, + status = hsa_executable_symbol_get_info(f._hsaSymbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, &privateSegmentSize); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); } - privateSegmentSize += sharedMemBytes; /* Kernel argument preparation. */ grid_launch_parm lp; - hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f->_kernelName); + hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel"); #if USE_DISPATCH_HSA_KERNEL @@ -315,7 +317,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, aql.grid_size_z = blockDimZ * gridDimZ; aql.group_segment_size = groupSegmentSize; aql.private_segment_size = privateSegmentSize; - aql.kernel_object = f->_kernel; + aql.kernel_object = f._object; aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | (1 << HSA_PACKET_HEADER_BARRIER) | @@ -356,7 +358,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, Launch AQL packet */ hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ, - gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_kernel); + gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_object); /* @@ -384,7 +386,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, #endif // USE_DISPATCH_HSA_KERNEL - ihipPostLaunchKernel(f->_kernelName, hStream, lp); + ihipPostLaunchKernel("ModuleKernel", hStream, lp); } @@ -405,9 +407,9 @@ hipError_t hipModuleGetGlobal(hipDeviceptr_t *dptr, size_t *bytes, } else{ hipFunction_t func; - ihipModuleGetFunction(&func, hmod, name); + ihipModuleGetSymbol(&func, hmod, name); *bytes = PrintSymbolSizes(hmod->ptr, name) + sizeof(amd_kernel_code_t); - *dptr = reinterpret_cast(func->_kernel); + *dptr = reinterpret_cast(func._object); return ihipLogStatus(ret); } } @@ -419,7 +421,7 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image) hipError_t ret = hipSuccess; if(image == NULL || module == NULL){ return ihipLogStatus(hipErrorNotInitialized); - }else{ + } else { auto ctx = ihipGetTlsDefaultCtx(); *module = new ihipModule_t; int deviceId = ctx->getDevice()->_deviceId; @@ -452,9 +454,13 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image) } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); + CHECK_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_freeze((*module)->executable, NULL); + CHECK_HSA(status, hipErrorNotInitialized); } return ihipLogStatus(ret); } diff --git a/hipamd/src/trace_helper.h b/hipamd/src/trace_helper.h index e75b492e0c..bde40d0690 100644 --- a/hipamd/src/trace_helper.h +++ b/hipamd/src/trace_helper.h @@ -72,6 +72,15 @@ inline std::string ToString(hipEvent_t v) return ss.str(); }; +// hipEvent_t specialization. TODO - maybe add an event ID for debug? +template <> +inline std::string ToString(hipFunction_t v) +{ + std::ostringstream ss; + ss << "0x" << std::hex << v._object; + return ss.str(); +}; + // hipStream_t From 06d382bc6d778edd782fc6583e3a807ad71fc018 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 17 Dec 2016 07:20:30 -0600 Subject: [PATCH 09/12] Remove USE_DISPATCH_HSA_KERNEL=0 path. --- hipamd/src/hip_hcc.cpp | 51 --------------------------------- hipamd/src/hip_hcc.h | 2 -- hipamd/src/hip_module.cpp | 60 --------------------------------------- 3 files changed, 113 deletions(-) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index 84794b02fb..b4568a8457 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -387,57 +387,6 @@ void ihipStream_t::lockclose_postKernelCommand(const char * kernelName, hc::acce }; - - -#if USE_DISPATCH_HSA_KERNEL==0 -// Precursor: the stream is already locked,specifically so this routine can enqueue work into the specified av. -void ihipStream_t::launchModuleKernel( - hc::accelerator_view av, - hsa_signal_t signal, - uint32_t blockDimX, - uint32_t blockDimY, - uint32_t blockDimZ, - uint32_t gridDimX, - uint32_t gridDimY, - uint32_t gridDimZ, - uint32_t groupSegmentSize, - uint32_t privateSegmentSize, - void *kernarg, - size_t kernSize, - uint64_t kernel){ - hsa_status_t status; - hsa_queue_t *Queue = (hsa_queue_t*)av.get_hsa_queue(); - const uint32_t queue_mask = Queue->size-1; - uint32_t packet_index = hsa_queue_load_write_index_relaxed(Queue); - hsa_kernel_dispatch_packet_t *dispatch_packet = &(((hsa_kernel_dispatch_packet_t*)(Queue->base_address))[packet_index & queue_mask]); - - dispatch_packet->completion_signal = signal; - dispatch_packet->workgroup_size_x = blockDimX; - dispatch_packet->workgroup_size_y = blockDimY; - dispatch_packet->workgroup_size_z = blockDimZ; - dispatch_packet->grid_size_x = blockDimX * gridDimX; - dispatch_packet->grid_size_y = blockDimY * gridDimY; - dispatch_packet->grid_size_z = blockDimZ * gridDimZ; - dispatch_packet->group_segment_size = groupSegmentSize; - dispatch_packet->private_segment_size = privateSegmentSize; - dispatch_packet->kernarg_address = kernarg; - dispatch_packet->kernel_object = kernel; - uint16_t header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | - (1 << HSA_PACKET_HEADER_BARRIER) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); - - uint16_t setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; - uint32_t header32 = header | (setup << 16); - - __atomic_store_n((uint32_t*)(dispatch_packet), header32, __ATOMIC_RELEASE); - - hsa_queue_store_write_index_relaxed(Queue, packet_index + 1); - hsa_signal_store_relaxed(Queue->doorbell_signal, packet_index); -} -#endif - - //============================================================================= // Recompute the peercnt and the packed _peerAgents whenever a peer is added or deleted. // The packed _peerAgents can efficiently be used on each memory allocation. diff --git a/hipamd/src/hip_hcc.h b/hipamd/src/hip_hcc.h index 4075c2b068..ac4d493862 100644 --- a/hipamd/src/hip_hcc.h +++ b/hipamd/src/hip_hcc.h @@ -33,9 +33,7 @@ THE SOFTWARE. #error("This version of HIP requires a newer version of HCC."); #endif -#define USE_DISPATCH_HSA_KERNEL 1 #define USE_IPC 0 -// //--- diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index cb56fb3b4a..65eab494b8 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -300,7 +300,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, grid_launch_parm lp; hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel"); -#if USE_DISPATCH_HSA_KERNEL hsa_kernel_dispatch_packet_t aql; @@ -325,65 +324,6 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize); -#else - - /* - Create signal - */ - - hsa_signal_t signal; - status = hsa_signal_create(1, 0, NULL, &signal); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } - - /* - Allocate kernarg - */ - void *kern = nullptr; - - hsa_amd_memory_pool_t *pool = reinterpret_cast(lp.av->get_hsa_kernarg_region()); - status = hsa_amd_memory_pool_allocate(*pool, kernArgSize, 0, &kern); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } - status = hsa_amd_agents_allow_access(1, (hsa_agent_t*)lp.av->get_hsa_agent(), 0, kern); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } - memcpy(kern, config[1], kernArgSize); - - - /* - Launch AQL packet - */ - hStream->launchModuleKernel(*lp.av, signal, blockDimX, blockDimY, blockDimZ, - gridDimX, gridDimY, gridDimZ, groupSegmentSize, privateSegmentSize, kern, kernArgSize, f->_object); - - - /* - Wait for signal - */ - - hsa_signal_value_t value = hsa_signal_wait_acquire(signal, HSA_SIGNAL_CONDITION_LT, 1, UINT64_MAX, HSA_WAIT_STATE_BLOCKED); - - /* - Destroy kernarg - */ - status = hsa_amd_memory_pool_free(kern); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } - - /* - Destroy the signal - */ - status = hsa_signal_destroy(signal); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } - -#endif // USE_DISPATCH_HSA_KERNEL ihipPostLaunchKernel("ModuleKernel", hStream, lp); From 2bd70ff3453a7b70f99874845c69f6dd47f29289 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 17 Dec 2016 07:21:15 -0600 Subject: [PATCH 10/12] Remove HSA dependency from hipFunction_t Place _groupSegmentSize and _privateSegmentSize inside Function, remove hsa_executable_symbol_t. --- .../include/hip/hcc_detail/hip_runtime_api.h | 8 ++- hipamd/src/hip_module.cpp | 63 ++++++++++--------- 2 files changed, 37 insertions(+), 34 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index e38f1d6982..0be4961f03 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -30,10 +30,10 @@ THE SOFTWARE. #include #include +#include #include #include -#include #if defined (__HCC__) && (__hcc_workweek__ < 16155) #error("This version of HIP requires a newer version of HCC."); @@ -73,8 +73,10 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t; typedef struct ihipModule_t *hipModule_t; struct ihipModuleSymbol_t{ - hsa_executable_symbol_t _hsaSymbol; - uint64_t _object; // The kernel object. + uint64_t _object; // The kernel object. + uint32_t _groupSegmentSize; + uint32_t _privateSegmentSize; + //std::string _name; // TODO - review for performance cost. Name is just used for debug. }; typedef struct ihipModuleSymbol_t hipFunction_t; diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 65eab494b8..19de6957e5 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -36,6 +36,11 @@ THE SOFTWARE. //TODO Use Pool APIs from HCC to get memory regions. #define CHECK_HSA(hsaStatus, hipStatus) \ +if (hsaStatus != HSA_STATUS_SUCCESS) {\ + return hipStatus;\ +} + +#define CHECKLOG_HSA(hsaStatus, hipStatus) \ if (hsaStatus != HSA_STATUS_SUCCESS) {\ return ihipLogStatus(hipStatus);\ } @@ -108,6 +113,7 @@ uint64_t ElfSize(const void *emi){ return total_size; } + hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ HIP_INIT_API(module, fname); hipError_t ret = hipSuccess; @@ -159,19 +165,20 @@ hipError_t hipModuleLoad(hipModule_t *module, const char *fname){ } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - CHECK_HSA(status, hipErrorNotInitialized); + CHECKLOG_HSA(status, hipErrorNotInitialized); status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); - CHECK_HSA(status, hipErrorNotInitialized); + CHECKLOG_HSA(status, hipErrorNotInitialized); status = hsa_executable_freeze((*module)->executable, NULL); - CHECK_HSA(status, hipErrorNotInitialized); + CHECKLOG_HSA(status, hipErrorNotInitialized); } } return ihipLogStatus(ret); } + hipError_t hipModuleUnload(hipModule_t hmod) { // TODO - improve this synchronization so it is thread-safe. @@ -214,21 +221,29 @@ hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char ihipDevice_t *currentDevice = ihipGetDevice(deviceId); hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent; - - hsa_status_t status; - status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &(func->_hsaSymbol)); + hsa_executable_symbol_t symbol; + status = hsa_executable_get_symbol(hmod->executable, NULL, name, gpuAgent, 0, &symbol); if(status != HSA_STATUS_SUCCESS){ return ihipLogStatus(hipErrorNotFound); } - status = hsa_executable_symbol_get_info(func->_hsaSymbol, + status = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, &func->_object); + CHECK_HSA(status, hipErrorNotFound); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } + status = hsa_executable_symbol_get_info(symbol, + HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, + &func->_groupSegmentSize); + CHECK_HSA(status, hipErrorNotFound); + + status = hsa_executable_symbol_get_info(symbol, + HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, + &func->_privateSegmentSize); + CHECK_HSA(status, hipErrorNotFound); + + //func->_name = std::string(name); } return ihipLogStatus(ret); } @@ -277,27 +292,13 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, return ihipLogStatus(hipErrorInvalidValue); } - uint32_t groupSegmentSize; - hsa_status_t status = hsa_executable_symbol_get_info(f._hsaSymbol, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_GROUP_SEGMENT_SIZE, - &groupSegmentSize); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } - - uint32_t privateSegmentSize; - status = hsa_executable_symbol_get_info(f._hsaSymbol, - HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_PRIVATE_SEGMENT_SIZE, - &privateSegmentSize); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotFound); - } /* Kernel argument preparation. */ grid_launch_parm lp; + //hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f._name.empty() ? "ModuleKernel" : f._name.c_str()); hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel"); @@ -314,8 +315,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, aql.grid_size_x = blockDimX * gridDimX; aql.grid_size_y = blockDimY * gridDimY; aql.grid_size_z = blockDimZ * gridDimZ; - aql.group_segment_size = groupSegmentSize; - aql.private_segment_size = privateSegmentSize; + aql.group_segment_size = f._groupSegmentSize + sharedMemBytes; + aql.private_segment_size = f._privateSegmentSize; aql.kernel_object = f._object; aql.setup = 3 << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; aql.header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | @@ -326,8 +327,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize); + //ihipPostLaunchKernel(f._name.empty() ? "ModuleKernel" : f._name.c_str(), hStream, lp); ihipPostLaunchKernel("ModuleKernel", hStream, lp); - } return ihipLogStatus(ret); @@ -394,13 +395,13 @@ hipError_t hipModuleLoadData(hipModule_t *module, const void *image) } status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &(*module)->executable); - CHECK_HSA(status, hipErrorNotInitialized); + CHECKLOG_HSA(status, hipErrorNotInitialized); status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); - CHECK_HSA(status, hipErrorNotInitialized); + CHECKLOG_HSA(status, hipErrorNotInitialized); status = hsa_executable_freeze((*module)->executable, NULL); - CHECK_HSA(status, hipErrorNotInitialized); + CHECKLOG_HSA(status, hipErrorNotInitialized); } return ihipLogStatus(ret); } From 5d815937de850d959d1096654abbd5d1f1375817 Mon Sep 17 00:00:00 2001 From: Ben Sander Date: Sat, 17 Dec 2016 08:54:09 -0600 Subject: [PATCH 11/12] Add name for function --- hipamd/include/hip/hcc_detail/hip_runtime_api.h | 3 +-- hipamd/src/hip_module.cpp | 10 ++++------ 2 files changed, 5 insertions(+), 8 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 0be4961f03..4c9cedecc4 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -30,7 +30,6 @@ THE SOFTWARE. #include #include -#include #include #include @@ -76,7 +75,7 @@ struct ihipModuleSymbol_t{ uint64_t _object; // The kernel object. uint32_t _groupSegmentSize; uint32_t _privateSegmentSize; - //std::string _name; // TODO - review for performance cost. Name is just used for debug. + char _name[64]; // TODO - review for performance cost. Name is just used for debug. }; typedef struct ihipModuleSymbol_t hipFunction_t; diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index 19de6957e5..74cdd8a4ae 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -243,7 +243,7 @@ hipError_t ihipModuleGetSymbol(hipFunction_t *func, hipModule_t hmod, const char &func->_privateSegmentSize); CHECK_HSA(status, hipErrorNotFound); - //func->_name = std::string(name); + strncpy(func->_name, name, sizeof(func->_name)); } return ihipLogStatus(ret); } @@ -297,9 +297,8 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, /* Kernel argument preparation. */ - grid_launch_parm lp; - //hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f._name.empty() ? "ModuleKernel" : f._name.c_str()); - hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, "ModuleKernel"); + grid_launch_parm lp; // TODO - dummy arg but values are printed during debug. + hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f._name); hsa_kernel_dispatch_packet_t aql; @@ -327,8 +326,7 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, lp.av->dispatch_hsa_kernel(&aql, config[1] /* kernarg*/, kernArgSize); - //ihipPostLaunchKernel(f._name.empty() ? "ModuleKernel" : f._name.c_str(), hStream, lp); - ihipPostLaunchKernel("ModuleKernel", hStream, lp); + ihipPostLaunchKernel(f._name, hStream, lp); } return ihipLogStatus(ret); From 4988975b5991b53c1baf1cd1e3f626a246210b92 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Wed, 23 Nov 2016 09:50:33 +0530 Subject: [PATCH 12/12] Fix for HCSWAP-67 Change-Id: I0b2ce5ab933237947fb41d89769db3da16e5be6a Conflicts: src/hip_hcc.cpp --- hipamd/src/hip_hcc.cpp | 11 ++++------- 1 file changed, 4 insertions(+), 7 deletions(-) diff --git a/hipamd/src/hip_hcc.cpp b/hipamd/src/hip_hcc.cpp index b4568a8457..19d7e1169d 100644 --- a/hipamd/src/hip_hcc.cpp +++ b/hipamd/src/hip_hcc.cpp @@ -732,13 +732,10 @@ hipError_t ihipDevice_t::initProperties(hipDeviceProp_t* prop) _isLargeBar = _acc.has_cpu_accessible_am(); // Get Max Threads Per Multiprocessor - - HsaNodeProperties node_prop = {0}; - if(HSAKMT_STATUS_SUCCESS == hsaKmtGetNodeProperties(node, &node_prop)) { - uint32_t waves_per_cu = node_prop.MaxWavesPerSIMD; - uint32_t simd_per_cu = node_prop.NumSIMDPerCU; - prop-> maxThreadsPerMultiProcessor = prop->warpSize*waves_per_cu*simd_per_cu; - } + uint32_t max_waves_per_cu; + err = hsa_agent_get_info(_hsaAgent,(hsa_agent_info_t) HSA_AMD_AGENT_INFO_MAX_WAVES_PER_CU, &max_waves_per_cu); + DeviceErrorCheck(err); + prop-> maxThreadsPerMultiProcessor = prop->warpSize*max_waves_per_cu; // Get memory properties err = hsa_agent_iterate_regions(_hsaAgent, get_region_info, prop);