From b863c1392c727e454b2b2a3166077a0b18779fad Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 21 Aug 2019 21:36:53 +0300 Subject: [PATCH 1/8] [HIPIFY][fix] Set the correct exit value for hipify-clang if errors. This fixes a hanging of unit tests. --- hipify-clang/src/main.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/hipify-clang/src/main.cpp b/hipify-clang/src/main.cpp index 702a02d234..3b8f454fd7 100644 --- a/hipify-clang/src/main.cpp +++ b/hipify-clang/src/main.cpp @@ -426,6 +426,7 @@ int main(int argc, const char **argv) { if (Tool.runAndSave(&actionFactory)) { currentStat.hasErrors = true; LLVM_DEBUG(llvm::dbgs() << "Skipped some replacements.\n"); + Result = 1; } // Copy the tmpfile to the output if (!NoOutput && !currentStat.hasErrors) { From 71559200c02cb3f63d1712d19e974fc94dac0e79 Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 23 Aug 2019 02:19:18 -0700 Subject: [PATCH 2/8] Fix memcpy with IPC slowness (#1321) * Fix memcpy with IPC slowness * Make early erroneous returns * Real Clean up * Real Clean up++ --- src/hip_memory.cpp | 72 +++++++++++++++++++++++++++++----------------- 1 file changed, 45 insertions(+), 27 deletions(-) diff --git a/src/hip_memory.cpp b/src/hip_memory.cpp index 66f776e966..a1455b14cd 100644 --- a/src/hip_memory.cpp +++ b/src/hip_memory.cpp @@ -2156,46 +2156,64 @@ hipError_t hipIpcGetMemHandle(hipIpcMemHandle_t* handle, void* devPtr) { hipError_t hipIpcOpenMemHandle(void** devPtr, hipIpcMemHandle_t handle, unsigned int flags) { HIP_INIT_API(hipIpcOpenMemHandle, devPtr, &handle, flags); hipError_t hipStatus = hipSuccess; - if (devPtr == NULL) { - hipStatus = hipErrorInvalidValue; - } else { -#if USE_IPC - // Get the current device agent. - hc::accelerator acc; - hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); - if (!agent) return hipErrorInvalidResourceHandle; + if (devPtr == NULL) + return ihipLogStatus(hipErrorInvalidValue); - ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*)&handle; - // Attach ipc memory - auto ctx = ihipGetTlsDefaultCtx(); - { - LockedAccessor_CtxCrit_t crit(ctx->criticalData()); - // the peerCnt always stores self so make sure the trace actually - hsa_status_t hsa_status = hsa_amd_ipc_memory_attach( - (hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, crit->peerCnt(), - crit->peerAgents(), devPtr); - if (hsa_status != HSA_STATUS_SUCCESS) hipStatus = hipErrorMapBufferObjectFailed; - } +#if USE_IPC + // Get the current device agent. + hc::accelerator acc; + hsa_agent_t* agent = static_cast(acc.get_hsa_agent()); + if (!agent) + return ihipLogStatus(hipErrorInvalidResourceHandle); + + ihipIpcMemHandle_t* iHandle = (ihipIpcMemHandle_t*)&handle; + // Attach ipc memory + auto ctx = ihipGetTlsDefaultCtx(); + { + LockedAccessor_CtxCrit_t crit(ctx->criticalData()); + auto device = ctx->getWriteableDevice(); + // the peerCnt always stores self so make sure the trace actually + if(hsa_amd_ipc_memory_attach( + (hsa_amd_ipc_memory_t*)&(iHandle->ipc_handle), iHandle->psize, crit->peerCnt(), + crit->peerAgents(), devPtr) != HSA_STATUS_SUCCESS) + return ihipLogStatus(hipErrorRuntimeOther); + + hc::AmPointerInfo ampi(NULL, *devPtr, *devPtr, sizeof(*devPtr), acc, true, true); + am_status_t am_status = hc::am_memtracker_add(*devPtr,ampi); + if (am_status != AM_SUCCESS) + return ihipLogStatus(hipErrorMapBufferObjectFailed); + +#if USE_APP_PTR_FOR_CTX + am_status = hc::am_memtracker_update(*devPtr, device->_deviceId, 0, ctx); #else - hipStatus = hipErrorRuntimeOther; + am_status = hc::am_memtracker_update(*devPtr, device->_deviceId, 0); #endif + if(am_status != AM_SUCCESS) + return ihipLogStatus(hipErrorMapBufferObjectFailed); } +#else + hipStatus = hipErrorRuntimeOther; +#endif + return ihipLogStatus(hipStatus); } hipError_t hipIpcCloseMemHandle(void* devPtr) { HIP_INIT_API(hipIpcCloseMemHandle, devPtr); hipError_t hipStatus = hipSuccess; - if (devPtr == NULL) { - hipStatus = hipErrorInvalidValue; - } else { + if (devPtr == NULL) + return ihipLogStatus(hipErrorInvalidValue); + #if USE_IPC - hsa_status_t hsa_status = hsa_amd_ipc_memory_detach(devPtr); - if (hsa_status != HSA_STATUS_SUCCESS) return hipErrorInvalidResourceHandle; + if(hc::am_memtracker_remove(devPtr) != AM_SUCCESS) + return ihipLogStatus(hipErrorInvalidValue); + + if (hsa_amd_ipc_memory_detach(devPtr) != HSA_STATUS_SUCCESS) + return ihipLogStatus(hipErrorInvalidResourceHandle); #else - hipStatus = hipErrorRuntimeOther; + hipStatus = hipErrorRuntimeOther; #endif - } + return ihipLogStatus(hipStatus); } From 5066700ace5eef442326607541257632bcafb421 Mon Sep 17 00:00:00 2001 From: Aryan Salmanpour Date: Fri, 23 Aug 2019 05:19:35 -0400 Subject: [PATCH 3/8] [hip] add initial implementation for hipLaunchCooperativeKernel API (#1339) * [hip] add initial implementation for hipLaunchCooperativeKernel API * [hip] use total number of work groups to initialize the GWS resource * [hip] use only one argument for init_gws kernel * [hip] use the device associated with the stream for checking the device properties --- include/hip/hcc_detail/device_library_decls.h | 4 + src/hip_module.cpp | 94 +++++++++++++++++++ 2 files changed, 98 insertions(+) diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 4e9772c80a..9a7636fc4a 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -59,6 +59,10 @@ extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float); extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float); extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float); +extern "C" __device__ __attribute__((convergent)) void __ockl_gws_init(uint nwm1, uint rid); +extern "C" __device__ __attribute__((convergent)) void __ockl_gws_barrier(uint nwm1, uint rid); + + // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/src/hip_module.cpp b/src/hip_module.cpp index 79e07161ac..50b3cc9b04 100644 --- a/src/hip_module.cpp +++ b/src/hip_module.cpp @@ -386,6 +386,100 @@ hipError_t hipExtLaunchMultiKernelMultiDevice(hipLaunchParams* launchParamsList, return ihipLogStatus(result); } +namespace { +// kernel for initializing GWS +// nwm1 is the total number of work groups minus 1 +__global__ void init_gws(uint nwm1) { + __ockl_gws_init(nwm1, 0); +} +} + +hipError_t hipLaunchCooperativeKernel(const void* f, dim3 gridDim, + dim3 blockDimX, void** kernelParams, unsigned int sharedMemBytes, + hipStream_t stream) { + + HIP_INIT_API(hipLaunchCooperativeKernel, f, gridDim, blockDimX, kernelParams, sharedMemBytes, stream); + hipError_t result; + + + if ((f == nullptr) || (stream == nullptr) || (kernelParams == nullptr)) { + return ihipLogStatus(hipErrorNotInitialized); + } + + if (!stream->getDevice()->_props.cooperativeLaunch) { + return ihipLogStatus(hipErrorInvalidConfiguration); + } + + // Prepare the kernel descriptor for initializing the GWS + hipFunction_t gwsKD = hip_impl::get_program_state().kernel_descriptor( + reinterpret_cast(&init_gws), + hip_impl::target_agent(stream)); + + if (gwsKD == nullptr) { + return ihipLogStatus(hipErrorInvalidValue); + } + hip_impl::kernargs_size_align gwsKargs = + hip_impl::get_program_state().get_kernargs_size_align( + reinterpret_cast(&init_gws)); + + gwsKD->_kernarg_layout = *reinterpret_cast>*>(gwsKargs.getHandle()); + + // Prepare the kernel descriptor for the main kernel + hipFunction_t kd = hip_impl::get_program_state().kernel_descriptor( + reinterpret_cast(f), + hip_impl::target_agent(stream)); + if (kd == nullptr) { + return ihipLogStatus(hipErrorInvalidValue); + } + hip_impl::kernargs_size_align kargs = + hip_impl::get_program_state().get_kernargs_size_align( + reinterpret_cast(f)); + + kd->_kernarg_layout = *reinterpret_cast>*>(kargs.getHandle()); + + + void *gwsKernelParam[1]; + // calculate total number of work groups minus 1 for the main kernel + uint nwm1 = (gridDim.x * gridDim.y * gridDim.z) - 1; + gwsKernelParam[0] = &nwm1; + + LockedAccessor_StreamCrit_t streamCrit(stream->criticalData(), false); +#if (__hcc_workweek__ >= 19213) + streamCrit->_av.acquire_locked_hsa_queue(); +#endif + + // launch the init_gws kernel to initialize the GWS + result = ihipModuleLaunchKernel(tls, gwsKD, 1, 1, 1, 1, 1, 1, + 0, stream, gwsKernelParam, nullptr, nullptr, nullptr, 0, true); + + if (result != hipSuccess) { + stream->criticalData().unlock(); +#if (__hcc_workweek__ >= 19213) + stream->criticalData()._av.release_locked_hsa_queue(); +#endif + + return ihipLogStatus(hipErrorLaunchFailure); + } + + // launch the main kernel + result = ihipModuleLaunchKernel(tls, kd, + gridDim.x * blockDimX.x, + gridDim.y * blockDimX.y, + gridDim.z * blockDimX.z, + blockDimX.x, blockDimX.y, blockDimX.z, + sharedMemBytes, stream, kernelParams, nullptr, nullptr, + nullptr, 0, true); + + stream->criticalData().unlock(); +#if (__hcc_workweek__ >= 19213) + stream->criticalData()._av.release_locked_hsa_queue(); +#endif + + return ihipLogStatus(result); +} + namespace hip_impl { hsa_executable_t executable_for(hipModule_t hmod) { return hmod->executable; From 344d150bf854bab1be6e072661dae753379bb64d Mon Sep 17 00:00:00 2001 From: chrispaquot <35546540+chrispaquot@users.noreply.github.com> Date: Fri, 23 Aug 2019 02:19:50 -0700 Subject: [PATCH 4/8] Set device before processing each one (#1358) --- samples/1_Utils/hipInfo/hipInfo.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index 23e014ad36..14455894c5 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/samples/1_Utils/hipInfo/hipInfo.cpp @@ -192,6 +192,7 @@ int main(int argc, char* argv[]) { HIPCHECK(hipGetDeviceCount(&deviceCnt)); for (int i = 0; i < deviceCnt; i++) { + hipSetDevice(i); printDeviceProp(i); } From 0fd14a3e1342adc94b4a83fb91536dca3a2142bd Mon Sep 17 00:00:00 2001 From: Rahul Garg Date: Fri, 23 Aug 2019 02:20:02 -0700 Subject: [PATCH 5/8] Make Bundled_code_header visible for hipRTC usage (#1359) --- include/hip/hcc_detail/code_object_bundle.hpp | 8 ++++++++ 1 file changed, 8 insertions(+) diff --git a/include/hip/hcc_detail/code_object_bundle.hpp b/include/hip/hcc_detail/code_object_bundle.hpp index 8b1ae8aa49..32b0c0dbc8 100644 --- a/include/hip/hcc_detail/code_object_bundle.hpp +++ b/include/hip/hcc_detail/code_object_bundle.hpp @@ -91,6 +91,10 @@ struct Bundled_code { #define magic_string_ "__CLANG_OFFLOAD_BUNDLE__" +#ifdef __GNUC__ +#pragma GCC visibility push (default) +#endif + class Bundled_code_header { // DATA - STATICS static constexpr auto magic_string_sz_ = sizeof(magic_string_) - 1; @@ -172,6 +176,10 @@ class Bundled_code_header { size_t bundled_code_size = 0; }; +#ifdef __GNUC__ +#pragma GCC visibility pop +#endif + // CREATORS template Bundled_code_header::Bundled_code_header(RandomAccessIterator f, RandomAccessIterator l) From 7257f230634a95a86f323764e2ca1ac472a45bcf Mon Sep 17 00:00:00 2001 From: Sarbojit2019 <52527887+SarbojitAMD@users.noreply.github.com> Date: Fri, 23 Aug 2019 14:51:47 +0530 Subject: [PATCH 6/8] Added missing device prop fields into hipInfo sample (#1357) * Added prop.integrated into hipInfo sample * Added missing deviceProp fileds in hipInfo --- samples/1_Utils/hipInfo/hipInfo.cpp | 13 ++++++++++++- 1 file changed, 12 insertions(+), 1 deletion(-) diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index 14455894c5..e17f19675a 100644 --- a/samples/1_Utils/hipInfo/hipInfo.cpp +++ b/samples/1_Utils/hipInfo/hipInfo.cpp @@ -83,6 +83,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "Name: " << props.name << endl; cout << setw(w1) << "pciBusID: " << props.pciBusID << endl; cout << setw(w1) << "pciDeviceID: " << props.pciDeviceID << endl; + cout << setw(w1) << "pciDomainID: " << props.pciDomainID << endl; cout << setw(w1) << "multiProcessorCount: " << props.multiProcessorCount << endl; cout << setw(w1) << "maxThreadsPerMultiProcessor: " << props.maxThreadsPerMultiProcessor << endl; @@ -100,6 +101,7 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "totalConstMem: " << props.totalConstMem << endl; cout << setw(w1) << "sharedMemPerBlock: " << (float)props.sharedMemPerBlock / 1024.0 << " KB" << endl; + cout << setw(w1) << "canMapHostMemory: " << props.canMapHostMemory << endl; cout << setw(w1) << "regsPerBlock: " << props.regsPerBlock << endl; cout << setw(w1) << "warpSize: " << props.warpSize << endl; cout << setw(w1) << "l2CacheSize: " << props.l2CacheSize << endl; @@ -114,6 +116,8 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "major: " << props.major << endl; cout << setw(w1) << "minor: " << props.minor << endl; cout << setw(w1) << "concurrentKernels: " << props.concurrentKernels << endl; + cout << setw(w1) << "cooperativeLaunch: " << props.cooperativeLaunch << endl; + cout << setw(w1) << "cooperativeMultiDeviceLaunch: " << props.cooperativeMultiDeviceLaunch << endl; cout << setw(w1) << "arch.hasGlobalInt32Atomics: " << props.arch.hasGlobalInt32Atomics << endl; cout << setw(w1) << "arch.hasGlobalFloatAtomicExch: " << props.arch.hasGlobalFloatAtomicExch << endl; @@ -134,7 +138,14 @@ void printDeviceProp(int deviceId) { cout << setw(w1) << "arch.has3dGrid: " << props.arch.has3dGrid << endl; cout << setw(w1) << "arch.hasDynamicParallelism: " << props.arch.hasDynamicParallelism << endl; cout << setw(w1) << "gcnArch: " << props.gcnArch << endl; - + cout << setw(w1) << "isIntegrated: " << props.integrated << endl; + cout << setw(w1) << "maxTexture1D: " << props.maxTexture1D << endl; + cout << setw(w1) << "maxTexture2D.width: " << props.maxTexture2D[0] << endl; + cout << setw(w1) << "maxTexture2D.height: " << props.maxTexture2D[1] << endl; + cout << setw(w1) << "maxTexture3D.width: " << props.maxTexture3D[0] << endl; + cout << setw(w1) << "maxTexture3D.height: " << props.maxTexture3D[1] << endl; + cout << setw(w1) << "maxTexture3D.depth: " << props.maxTexture3D[2] << endl; + int deviceCnt; hipGetDeviceCount(&deviceCnt); cout << setw(w1) << "peers: "; From e1d4f8510aeca03e085aff8d12c1bb5f86331e03 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 23 Aug 2019 19:32:24 +0300 Subject: [PATCH 7/8] [HIPIFY][cmake] Explicitly set c++14 [Reason] LLVM became c++14 last week due to the following change: 37508d3dd94b0154861a90b1909d17b01400df99 Replace llvm::integer_sequence and friends with the C++14 standard version --- hipify-clang/CMakeLists.txt | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/hipify-clang/CMakeLists.txt b/hipify-clang/CMakeLists.txt index 6853cafaa1..7ba92cd93d 100644 --- a/hipify-clang/CMakeLists.txt +++ b/hipify-clang/CMakeLists.txt @@ -74,12 +74,12 @@ if ((LLVM_PACKAGE_VERSION VERSION_EQUAL "7") OR (LLVM_PACKAGE_VERSION VERSION_GR endif() set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") -set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS}") +set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} ${EXTRA_CFLAGS} -std=c++14") if(MSVC) target_compile_options(hipify-clang PRIVATE "/Od /GR- /EHs- /EHc-") set(CMAKE_CXX_LINK_FLAGS "${CMAKE_CXX_LINK_FLAGS} /SUBSYSTEM:WINDOWS") else() - set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -std=c++11 -pthread -fno-rtti -fvisibility-inlines-hidden") + set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -pthread -fno-rtti -fvisibility-inlines-hidden") endif() set(CMAKE_CXX_FLAGS "${CMAKE_CXX_FLAGS} -DHIPIFY_CLANG_RES=\\\"${LLVM_LIBRARY_DIRS}/clang/${LLVM_VERSION_MAJOR}.${LLVM_VERSION_MINOR}.${LLVM_VERSION_PATCH}\\\"") From 3a6ca29815751df6cea38a750d700078c4e5ea1d Mon Sep 17 00:00:00 2001 From: ramcherukuri <54456794+ramcherukuri@users.noreply.github.com> Date: Fri, 23 Aug 2019 16:36:33 -0700 Subject: [PATCH 8/8] moving result_of_t to result_of --- include/hip/hcc_detail/helpers.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/helpers.hpp b/include/hip/hcc_detail/helpers.hpp index 9366885f8b..e69cac4c21 100644 --- a/include/hip/hcc_detail/helpers.hpp +++ b/include/hip/hcc_detail/helpers.hpp @@ -89,7 +89,7 @@ template struct is_callable_impl : std::false_type {}; template -struct is_callable_impl > > : std::true_type {}; +struct is_callable_impl std::result_of::type(Ts...) > > : std::true_type {}; #else // C++17