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}\\\"") 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) { 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) 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/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 diff --git a/samples/1_Utils/hipInfo/hipInfo.cpp b/samples/1_Utils/hipInfo/hipInfo.cpp index 23e014ad36..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: "; @@ -192,6 +203,7 @@ int main(int argc, char* argv[]) { HIPCHECK(hipGetDeviceCount(&deviceCnt)); for (int i = 0; i < deviceCnt; i++) { + hipSetDevice(i); printDeviceProp(i); } 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); } 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;