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. 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()) { 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/include/hip/hcc_detail/hip_runtime_api.h b/hipamd/include/hip/hcc_detail/hip_runtime_api.h index 52e59ed17a..4c9cedecc4 100644 --- a/hipamd/include/hip/hcc_detail/hip_runtime_api.h +++ b/hipamd/include/hip/hcc_detail/hip_runtime_api.h @@ -33,7 +33,6 @@ THE SOFTWARE. #include #include -//#include "hip/hip_hcc.h" #if defined (__HCC__) && (__hcc_workweek__ < 16155) #error("This version of HIP requires a newer version of HCC."); @@ -53,7 +52,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; @@ -72,7 +71,14 @@ typedef struct ihipIpcEventHandle_t *hipIpcEventHandle_t; typedef struct ihipModule_t *hipModule_t; -typedef struct ihipFunction_t *hipFunction_t; +struct ihipModuleSymbol_t{ + uint64_t _object; // The kernel object. + uint32_t _groupSegmentSize; + uint32_t _privateSegmentSize; + char _name[64]; // TODO - review for performance cost. Name is just used for debug. +}; + +typedef struct ihipModuleSymbol_t hipFunction_t; typedef void* hipDeviceptr_t; @@ -1904,18 +1910,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/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: - - 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; 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. */ 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/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_hcc.cpp b/hipamd/src/hip_hcc.cpp index 84794b02fb..19d7e1169d 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. @@ -783,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); diff --git a/hipamd/src/hip_hcc.h b/hipamd/src/hip_hcc.h index b5417080e0..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 -// //--- @@ -383,26 +381,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 +390,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_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 } diff --git a/hipamd/src/hip_module.cpp b/hipamd/src/hip_module.cpp index badeac2dcd..74cdd8a4ae 100644 --- a/hipamd/src/hip_module.cpp +++ b/hipamd/src/hip_module.cpp @@ -35,6 +35,16 @@ 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);\ +} + namespace hipdrv { hsa_status_t findSystemRegions(hsa_region_t region, void *data){ @@ -103,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; @@ -154,15 +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); - if(status != HSA_STATUS_SUCCESS){ - return ihipLogStatus(hipErrorNotInitialized); - } + CHECKLOG_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); + CHECKLOG_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_freeze((*module)->executable, NULL); + CHECKLOG_HSA(status, hipErrorNotInitialized); } } return ihipLogStatus(ret); } + hipError_t hipModuleUnload(hipModule_t hmod) { // TODO - improve this synchronization so it is thread-safe. @@ -189,7 +205,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,31 +217,33 @@ 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_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)->_kernelSymbol, + status = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT, - &(*func)->_kernel); + &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); + + strncpy(func->_name, name, sizeof(func->_name)); } return ihipLogStatus(ret); } @@ -234,7 +252,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); } @@ -274,31 +292,14 @@ hipError_t hipModuleLaunchKernel(hipFunction_t f, return ihipLogStatus(hipErrorInvalidValue); } - uint32_t groupSegmentSize; - hsa_status_t status = hsa_executable_symbol_get_info(f->_kernelSymbol, - 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->_kernelSymbol, - 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); + grid_launch_parm lp; // TODO - dummy arg but values are printed during debug. + hStream = ihipPreLaunchKernel(hStream, 0, 0, &lp, f._name); -#if USE_DISPATCH_HSA_KERNEL hsa_kernel_dispatch_packet_t aql; @@ -313,9 +314,9 @@ 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.kernel_object = f->_kernel; + 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) | (1 << HSA_PACKET_HEADER_BARRIER) | @@ -323,69 +324,9 @@ 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->_kernel); - - - /* - 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(f->_kernelName, hStream, lp); - + ihipPostLaunchKernel(f._name, hStream, lp); } return ihipLogStatus(ret); @@ -405,9 +346,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 +360,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 +393,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); - } + CHECKLOG_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_load_code_object((*module)->executable, agent, (*module)->object, NULL); + CHECKLOG_HSA(status, hipErrorNotInitialized); + + status = hsa_executable_freeze((*module)->executable, NULL); + CHECKLOG_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 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;