From 1f7b924e2e7304202db3c7ff1f783b36c7417ef5 Mon Sep 17 00:00:00 2001 From: Yaxun Sam Liu Date: Sun, 21 Oct 2018 10:03:33 -0400 Subject: [PATCH 1/5] Let hipcc handle gfx906 for hip-clang --- hipamd/bin/hipcc | 6 +++++- 1 file changed, 5 insertions(+), 1 deletion(-) diff --git a/hipamd/bin/hipcc b/hipamd/bin/hipcc index f22884ea1b..5942b20b8d 100755 --- a/hipamd/bin/hipcc +++ b/hipamd/bin/hipcc @@ -687,7 +687,11 @@ if($HIP_PLATFORM eq "hcc" or $HIP_PLATFORM eq "clang"){ $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX900__=1 "; } if ($target_gfx906 eq 1) { - $HIPLDFLAGS .= " --amdgpu-target=gfx906"; + $GPU_ARCH_ARG = $GPU_ARCH_OPT . "gfx906"; + $HIPLDFLAGS .= $GPU_ARCH_ARG; + if ($HIP_PLATFORM eq 'clang') { + $HIPCXXFLAGS .= $GPU_ARCH_ARG;; + } $HIPCXXFLAGS .= " -D__HIP_ARCH_GFX906__=1 "; } } From b01ac269486a356cb05dc0c1d8e7a3480b30afaf Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Tue, 23 Oct 2018 13:20:49 +0530 Subject: [PATCH 2/5] Make HIP functional again with HCC from ROCm 1.9.x Change-Id: I214acdfd0b79dcf783993e44fe31baee64fd4dc3 --- hipamd/Jenkinsfile | 49 +++++++++++- .../include/hip/hcc_detail/device_functions.h | 74 +++++++++++++++++++ .../hip/hcc_detail/device_library_decls.h | 24 ++++++ hipamd/src/hip_stream.cpp | 17 +++++ 4 files changed, 162 insertions(+), 2 deletions(-) diff --git a/hipamd/Jenkinsfile b/hipamd/Jenkinsfile index 540e7e6689..6d37f10e3c 100644 --- a/hipamd/Jenkinsfile +++ b/hipamd/Jenkinsfile @@ -296,8 +296,53 @@ def docker_upload_dockerhub( String local_org, String image_name, String remote_ String build_config = 'Release' String job_name = env.JOB_NAME.toLowerCase( ) -// The following launches 2 builds in parallel: rocm-head and cuda-9.x -parallel rocm_head: +// The following launches 3 builds in parallel: rocm-head, rocm-1.9.x and cuda-9.x +parallel rocm_1_9: +{ + node('hip-rocm') + { + String hcc_ver = 'rocm-1.9.x' + String from_image = 'ci_test_nodes/rocm-1.9.x/ubuntu-16.04:latest' + String inside_args = '--device=/dev/kfd --device=/dev/dri --group-add=video' + + // Checkout source code, dependencies and version files + String source_hip_rel = checkout_and_version( hcc_ver ) + + // Create/reuse a docker image that represents the hip build environment + def hip_build_image = docker_build_image( hcc_ver, 'hip', '', source_hip_rel, from_image ) + + // Print system information for the log + hip_build_image.inside( inside_args ) + { + sh """#!/usr/bin/env bash + set -x + /opt/rocm/bin/rocm_agent_enumerator -t ALL + /opt/rocm/bin/hcc --version + """ + } + + // Conctruct a binary directory path based on build config + String build_hip_rel = build_directory_rel( build_config ); + + // Build hip inside of the build environment + docker_build_inside_image( hip_build_image, inside_args, hcc_ver, '', build_config, source_hip_rel, build_hip_rel ) + + // Clean docker build image + docker_clean_images( 'hip', docker_build_image_name( ) ) + + // After a successful build, upload a docker image of the results + /* + String hip_image_name = docker_upload_artifactory( hcc_ver, job_name, from_image, source_hip_rel, build_hip_rel ) + if( params.push_image_to_docker_hub ) + { + docker_upload_dockerhub( job_name, hip_image_name, 'rocm' ) + docker_clean_images( 'rocm', hip_image_name ) + } + docker_clean_images( job_name, hip_image_name ) + */ + } +}, +rocm_head: { node('hip-rocm') { diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 42927e3246..2d5f73c153 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -752,6 +752,80 @@ void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } +#if defined(__HCC__) && (__hcc_minor__ < 3) +// hip.amdgcn.bc - sync threads +#define __CLK_LOCAL_MEM_FENCE 0x01 +typedef unsigned __cl_mem_fence_flags; + +typedef enum __memory_scope { + __memory_scope_work_item = __OPENCL_MEMORY_SCOPE_WORK_ITEM, + __memory_scope_work_group = __OPENCL_MEMORY_SCOPE_WORK_GROUP, + __memory_scope_device = __OPENCL_MEMORY_SCOPE_DEVICE, + __memory_scope_all_svm_devices = __OPENCL_MEMORY_SCOPE_ALL_SVM_DEVICES, + __memory_scope_sub_group = __OPENCL_MEMORY_SCOPE_SUB_GROUP +} __memory_scope; + +// enum values aligned with what clang uses in EmitAtomicExpr() +typedef enum __memory_order +{ + __memory_order_relaxed = __ATOMIC_RELAXED, + __memory_order_acquire = __ATOMIC_ACQUIRE, + __memory_order_release = __ATOMIC_RELEASE, + __memory_order_acq_rel = __ATOMIC_ACQ_REL, + __memory_order_seq_cst = __ATOMIC_SEQ_CST +} __memory_order; + +__device__ +inline +static void +__atomic_work_item_fence(__cl_mem_fence_flags flags, __memory_order order, __memory_scope scope) +{ + // We're tying global-happens-before and local-happens-before together as does HSA + if (order != __memory_order_relaxed) { + switch (scope) { + case __memory_scope_work_item: + break; + case __memory_scope_sub_group: + switch (order) { + case __memory_order_relaxed: break; + case __memory_order_acquire: __llvm_fence_acq_sg(); break; + case __memory_order_release: __llvm_fence_rel_sg(); break; + case __memory_order_acq_rel: __llvm_fence_ar_sg(); break; + case __memory_order_seq_cst: __llvm_fence_sc_sg(); break; + } + break; + case __memory_scope_work_group: + switch (order) { + case __memory_order_relaxed: break; + case __memory_order_acquire: __llvm_fence_acq_wg(); break; + case __memory_order_release: __llvm_fence_rel_wg(); break; + case __memory_order_acq_rel: __llvm_fence_ar_wg(); break; + case __memory_order_seq_cst: __llvm_fence_sc_wg(); break; + } + break; + case __memory_scope_device: + switch (order) { + case __memory_order_relaxed: break; + case __memory_order_acquire: __llvm_fence_acq_dev(); break; + case __memory_order_release: __llvm_fence_rel_dev(); break; + case __memory_order_acq_rel: __llvm_fence_ar_dev(); break; + case __memory_order_seq_cst: __llvm_fence_sc_dev(); break; + } + break; + case __memory_scope_all_svm_devices: + switch (order) { + case __memory_order_relaxed: break; + case __memory_order_acquire: __llvm_fence_acq_sys(); break; + case __memory_order_release: __llvm_fence_rel_sys(); break; + case __memory_order_acq_rel: __llvm_fence_ar_sys(); break; + case __memory_order_seq_cst: __llvm_fence_sc_sys(); break; + } + break; + } + } +} +#endif + // Memory Fence Functions __device__ inline diff --git a/hipamd/include/hip/hcc_detail/device_library_decls.h b/hipamd/include/hip/hcc_detail/device_library_decls.h index 2a0b968fe3..2bf3c8cc51 100644 --- a/hipamd/include/hip/hcc_detail/device_library_decls.h +++ b/hipamd/include/hip/hcc_detail/device_library_decls.h @@ -65,6 +65,29 @@ extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float) __device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; } #endif //__HIP_DEVICE_COMPILE__ +#if defined(__HCC__) && (__hcc_minor__ < 3) +// __llvm_fence* functions from device-libs/irif/src/fence.ll +extern "C" __device__ void __llvm_fence_acq_sg(void); +extern "C" __device__ void __llvm_fence_acq_wg(void); +extern "C" __device__ void __llvm_fence_acq_dev(void); +extern "C" __device__ void __llvm_fence_acq_sys(void); + +extern "C" __device__ void __llvm_fence_rel_sg(void); +extern "C" __device__ void __llvm_fence_rel_wg(void); +extern "C" __device__ void __llvm_fence_rel_dev(void); +extern "C" __device__ void __llvm_fence_rel_sys(void); + +extern "C" __device__ void __llvm_fence_ar_sg(void); +extern "C" __device__ void __llvm_fence_ar_wg(void); +extern "C" __device__ void __llvm_fence_ar_dev(void); +extern "C" __device__ void __llvm_fence_ar_sys(void); + + +extern "C" __device__ void __llvm_fence_sc_sg(void); +extern "C" __device__ void __llvm_fence_sc_wg(void); +extern "C" __device__ void __llvm_fence_sc_dev(void); +extern "C" __device__ void __llvm_fence_sc_sys(void); +#else // Using hip.amdgcn.bc - sync threads #define __CLK_LOCAL_MEM_FENCE 0x01 typedef unsigned __cl_mem_fence_flags; @@ -90,5 +113,6 @@ typedef enum __memory_order // Linked from hip.amdgcn.bc extern "C" __device__ void __atomic_work_item_fence(__cl_mem_fence_flags, __memory_order, __memory_scope); +#endif #endif diff --git a/hipamd/src/hip_stream.cpp b/hipamd/src/hip_stream.cpp index 1e239cb2fc..2268a203dd 100644 --- a/hipamd/src/hip_stream.cpp +++ b/hipamd/src/hip_stream.cpp @@ -31,12 +31,21 @@ THE SOFTWARE. //------------------------------------------------------------------------------------------------- // Stream // +#if defined(__HCC__) && (__hcc_minor__ < 3) +enum queue_priority +{ + priority_high = 0, + priority_normal = 0, + priority_low = 0 +}; +#else enum queue_priority { priority_high = Kalmar::priority_high, priority_normal = Kalmar::priority_normal, priority_low = Kalmar::priority_low }; +#endif //--- hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priority) { @@ -59,7 +68,11 @@ hipError_t ihipStreamCreate(hipStream_t* stream, unsigned int flags, int priorit // Obtain mutex access to the device critical data, release by destructor LockedAccessor_CtxCrit_t ctxCrit(ctx->criticalData()); +#if defined(__HCC__) && (__hcc_minor__ < 3) + auto istream = new ihipStream_t(ctx, acc.create_view(), flags); +#else auto istream = new ihipStream_t(ctx, acc.create_view(Kalmar::execute_in_order, Kalmar::queuing_mode_automatic, (Kalmar::queue_priority)priority), flags); +#endif ctxCrit->addStream(istream); *stream = istream; @@ -223,8 +236,12 @@ hipError_t hipStreamGetPriority(hipStream_t stream, int* priority) { } else if (stream == hipStreamNull) { return ihipLogStatus(hipErrorInvalidResourceHandle); } else { +#if defined(__HCC__) && (__hcc_minor__ < 3) + *priority = 0; +#else LockedAccessor_StreamCrit_t crit(stream->_criticalData); *priority = crit->_av.get_queue_priority(); +#endif return ihipLogStatus(hipSuccess); } } From 7b6d23267232aa8d3d68938cddbabfa988640827 Mon Sep 17 00:00:00 2001 From: Maneesh Gupta Date: Wed, 24 Oct 2018 11:47:41 +0530 Subject: [PATCH 3/5] Fix logic for double variants of __shfl* Change-Id: I604f00b54cf4bd9c5f26ca6fa680fca5e9629417 --- .../include/hip/hcc_detail/device_functions.h | 52 ++++++++++++++----- 1 file changed, 40 insertions(+), 12 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/device_functions.h b/hipamd/include/hip/hcc_detail/device_functions.h index 42927e3246..f0b93895a0 100644 --- a/hipamd/include/hip/hcc_detail/device_functions.h +++ b/hipamd/include/hip/hcc_detail/device_functions.h @@ -297,9 +297,16 @@ float __shfl(float var, int src_lane, int width = warpSize) { __device__ inline double __shfl(double var, int src_lane, int width = warpSize) { - __u tmp; tmp.f = (float) var; - tmp.i = __shfl(tmp.i, src_lane, width); - return (double) tmp.f; + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl(tmp[0], src_lane, width); + tmp[1] = __shfl(tmp[1], src_lane, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; } __device__ @@ -327,9 +334,16 @@ float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) { __device__ inline double __shfl_up(double var, unsigned int lane_delta, int width = warpSize) { - __u tmp; tmp.f = (float) var; - tmp.i = __shfl_up(tmp.i, lane_delta, width); - return (double) tmp.f; + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_up(tmp[0], lane_delta, width); + tmp[1] = __shfl_up(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; } __device__ @@ -357,9 +371,16 @@ float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) { __device__ inline double __shfl_down(double var, unsigned int lane_delta, int width = warpSize) { - __u tmp; tmp.f = (float) var; - tmp.i = __shfl_down(tmp.i, lane_delta, width); - return (double) tmp.f; + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_down(tmp[0], lane_delta, width); + tmp[1] = __shfl_down(tmp[1], lane_delta, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; } __device__ @@ -387,9 +408,16 @@ float __shfl_xor(float var, int lane_mask, int width = warpSize) { __device__ inline double __shfl_xor(double var, int lane_mask, int width = warpSize) { - __u tmp; tmp.f = (float) var; - tmp.i = __shfl_xor(tmp.i, lane_mask, width); - return (double) tmp.f; + static_assert(sizeof(double) == 2 * sizeof(int), ""); + static_assert(sizeof(double) == sizeof(uint64_t), ""); + + int tmp[2]; __builtin_memcpy(tmp, &var, sizeof(tmp)); + tmp[0] = __shfl_xor(tmp[0], lane_mask, width); + tmp[1] = __shfl_xor(tmp[1], lane_mask, width); + + uint64_t tmp0 = (static_cast(tmp[1]) << 32ull) | static_cast(tmp[0]); + double tmp1; __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); + return tmp1; } #define MASK1 0x00ff00ff From b244d1fc99f7089f41a1d52d941184bbed536e83 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Thu, 25 Oct 2018 17:09:18 +0300 Subject: [PATCH 4/5] [HIPIFY][cmake][fix] Mark all target_link_libraries with PRIVATE keyword [Reason] Avoid cmake error occurred when mixing signatures with AddLLVM.cmake [Error] CMake Error at CMakeLists.txt:26 (target_link_libraries): The keyword signature for target_link_libraries has already been used with the target "hipify-clang". All uses of target_link_libraries with a target must be either all-keyword or all-plain. Error occurred only on Linux and LLVM 7.0.0 --- hipamd/hipify-clang/CMakeLists.txt | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/hipamd/hipify-clang/CMakeLists.txt b/hipamd/hipify-clang/CMakeLists.txt index a3952ba418..56de7c7904 100644 --- a/hipamd/hipify-clang/CMakeLists.txt +++ b/hipamd/hipify-clang/CMakeLists.txt @@ -23,7 +23,7 @@ set(CMAKE_CXX_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang++) set(CMAKE_C_COMPILER ${LLVM_TOOLS_BINARY_DIR}/clang) # Link against LLVM and CLANG libraries -target_link_libraries(hipify-clang +target_link_libraries(hipify-clang PRIVATE clangASTMatchers clangFrontend clangTooling @@ -48,11 +48,11 @@ target_link_libraries(hipify-clang LLVMCore) if(WIN32) - target_link_libraries(hipify-clang version) + target_link_libraries(hipify-clang PRIVATE version) endif() if ((LLVM_PACKAGE_VERSION VERSION_EQUAL "7") OR (LLVM_PACKAGE_VERSION VERSION_GREATER "7")) - target_link_libraries(hipify-clang clangToolingInclusions) + target_link_libraries(hipify-clang PRIVATE clangToolingInclusions) endif() set(CMAKE_C_FLAGS "${CMAKE_C_FLAGS} ${EXTRA_CFLAGS}") From 5ab9c4bbca426acbbfe7b7a9ef30590f0cb195f5 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 26 Oct 2018 15:01:58 +0300 Subject: [PATCH 5/5] [HIPIFY][doc] Readme update Clang 7.0.0 release has a known bug https://bugs.llvm.org/show_bug.cgi?id=36384, thus 7.0.0 is unstable release for Linux and Windows. Last stable and tested LLVM+Clang release is 6.0.1. --- hipamd/hipify-clang/README.md | 26 +++++++++++++------------- 1 file changed, 13 insertions(+), 13 deletions(-) diff --git a/hipamd/hipify-clang/README.md b/hipamd/hipify-clang/README.md index 02bc9cfed0..ec46c7cb2f 100644 --- a/hipamd/hipify-clang/README.md +++ b/hipamd/hipify-clang/README.md @@ -30,11 +30,7 @@ ## Dependencies `hipify-clang` requires: -1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended releases: - - 6.0.1 (linux and windows), - - 7.0.0 (linux only). +1. LLVM+CLANG of at least version 3.8.0, latest stable and recommended release: 6.0.1 (linux and windows). 2. CUDA at least version 7.5, latest supported release is 9.2. @@ -50,8 +46,8 @@ | 5.0.1 | 8.0 | | 5.0.2 | 8.0 | | 6.0.0 | 9.0 | -| 6.0.1 | 9.0 | -| 7.0.0 | 9.2 | linux only | +| **6.0.1** | **9.0** | **LATEST STABLE RELEASE** | +| 7.0.0 | 9.2 | windows is not supported, on linux there is a clang bug: https://bugs.llvm.org/show_bug.cgi?id=36384 | | | 10.0 | not yet supported | In most cases, you can get a suitable version of LLVM+CLANG with your package manager. @@ -77,10 +73,10 @@ cmake \ make -j install ``` -On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make`. +On Windows, the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make.` -Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested, `LLVM+CLANG` should be built in `Debug` mode as well. -64 bit build mode `-Thost=x64` is supported as well, `LLVM+CLANG` should be built (installed) in 64bit mode as well. +Debug build type `-DCMAKE_BUILD_TYPE=Debug` is also supported and tested; `LLVM+CLANG` should be built in `Debug` mode as well. +64 bit build mode `-Thost=x64` is supported as well; `LLVM+CLANG` should be built in 64bit mode as well. The binary can then be found at `./dist/bin/hipify-clang`. @@ -108,11 +104,11 @@ To run it: make -j install ``` - On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"` and after `cmake` the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`. + On Windows the following option should be specified for `cmake` at first place: `-G "Visual Studio 15 2017 Win64"`; the generated `LLVM.sln` should be built by `Visual Studio 15 2017` instead of `make`. 3. Ensure [`CUDA`](https://developer.nvidia.com/cuda-toolkit-archive) of minimum version 7.5 is installed. - * Having multiple CUDA installations in order to choose a concrete version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified: + * Having multiple CUDA installations, in order to choose a particular version the `DCUDA_TOOLKIT_ROOT_DIR` option should be specified: `-DCUDA_TOOLKIT_ROOT_DIR="C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0"` @@ -127,7 +123,9 @@ To run it: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-9.0-windows10-x64-v7.1` 5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed. + 6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM. + * installing `lit` into `python` might be required: `python f:/LLVM/6.0.1/llvm/utils/lit/setup.py install`, @@ -139,7 +137,9 @@ To run it: `-DLLVM_EXTERNAL_LIT=f:/LLVM/6.0.1/build/Release/bin/llvm-lit.py`, where `f:/LLVM/6.0.1/build/Release` is LLVM build directory. + 7. Build with the `HIPIFY_CLANG_TESTS` option turned on: -DHIPIFY_CLANG_TESTS=1. + 8. `make test-hipify` On Windows after `cmake` the project `test-hipify` in the generated `hipify-clang.sln` should be built by `Visual Studio 15 2017` instead of `make test-hipify`. @@ -178,7 +178,7 @@ A corresponding successful output: -- - CMake module path: F:/LLVM/6.0.1/dist/lib/cmake/llvm -- - Include path : F:/LLVM/6.0.1/dist/include -- - Binary path : F:/LLVM/6.0.1/dist/bin --- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "2.7") +-- Found PythonInterp: C:/Program Files/Python36/python.exe (found suitable version "3.6.4", minimum required is "3.6") -- Found lit: C:/Program Files/Python36/Scripts/lit.exe -- Found FileCheck: F:/LLVM/6.0.1/dist/bin/FileCheck.exe -- Found CUDA: C:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v9.0 (found version "9.0")