From cfe37484c9e35389a6a490b8f4efe7d5f1735690 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 14 Jun 2018 17:49:35 +0000 Subject: [PATCH] Split __llvm and device lib funcs into new headers --- include/hip/hcc_detail/device_library_decls.h | 51 +++++++++++++++++++ include/hip/hcc_detail/hip_runtime.h | 38 +++----------- include/hip/hcc_detail/llvm_intrinsics.h | 46 +++++++++++++++++ 3 files changed, 103 insertions(+), 32 deletions(-) create mode 100644 include/hip/hcc_detail/device_library_decls.h create mode 100644 include/hip/hcc_detail/llvm_intrinsics.h diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h new file mode 100644 index 0000000000..b3fa556bd4 --- /dev/null +++ b/include/hip/hcc_detail/device_library_decls.h @@ -0,0 +1,51 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +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, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @file hcc_detail/device_library_decls.h + * @brief Contains declarations for types and functions in device library. + */ + +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_LIBRARY_DECLS_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_LIBRARY_DECLS_H + +#include "hip/hcc_detail/host_defines.h" + +extern "C" __device__ int32_t __ockl_activelane_u32(void); + +// Introduce local address space +#define __local __attribute__((address_space(3))) +__device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } +extern "C" __device__ void* __local_to_generic(__local void* p); + +// __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); + +#endif diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 1589f19395..4cd41a0c86 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -62,6 +62,8 @@ THE SOFTWARE. #define CUDA_SUCCESS hipSuccess #include +#include +#include #endif // __HCC_OR_HIP_CLANG__ #if __HCC__ @@ -189,9 +191,6 @@ extern int HIP_TRACE_API; static constexpr int warpSize = 64; #define clock_t long long int -__device__ -unsigned long __llvm_amdgcn_s_memrealtime(void) __asm("llvm.amdgcn.s.memrealtime"); - __device__ inline long long int __clock64() { return (long long int)__llvm_amdgcn_s_memrealtime(); } @@ -221,8 +220,6 @@ int64_t __ballot64(int a) { } // hip.amdgcn.bc - lanemask -extern "C" __device__ int32_t __ockl_activelane_u32(void); - __device__ inline int64_t __lanemask_gt() @@ -285,27 +282,18 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask __host__ __device__ int min(int arg1, int arg2); __host__ __device__ int max(int arg1, int arg2); -// Introduce local address space -#define __local __attribute__((address_space(3))) -__device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } -extern "C" __device__ void* __local_to_generic(__local void* p); - -__device__ unsigned __llvm_amdgcn_s_getreg(unsigned) __asm("llvm.amdgcn.s.getreg"); - -__device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); __device__ inline void* __get_dynamicgroupbaseptr() { // Get group segment base pointer. - unsigned lds_base = __llvm_amdgcn_s_getreg(14342) << 8; - __local char* base = __to_local(lds_base); - unsigned long long group_static_size = __llvm_amdgcn_groupstaticsize(); - return (char*)__local_to_generic(base + group_static_size); + return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); } -__device__ inline void *__amdgcn_get_dynamicgroupbaseptr() { +__device__ +inline +void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } @@ -384,8 +372,6 @@ __device__ void __threadfence_system(void); */ // hip.amdgcn.bc - named sync -__device__ void __llvm_amdgcn_s_barrier() __asm("llvm.amdgcn.s.barrier"); - __device__ inline void __named_sync(int a, int b) { __llvm_amdgcn_s_barrier(); } #endif // __HCC_OR_HIP_CLANG__ @@ -669,7 +655,6 @@ void __assertfail(const char * __assertion, } // hip.amdgcn.bc - sync threads -// extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); #define __CLK_LOCAL_MEM_FENCE 0x01 typedef unsigned __cl_mem_fence_flags; @@ -691,17 +676,6 @@ typedef enum __memory_order __memory_order_seq_cst = __ATOMIC_SEQ_CST } __memory_order; -// __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); - __device__ inline static void __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) diff --git a/include/hip/hcc_detail/llvm_intrinsics.h b/include/hip/hcc_detail/llvm_intrinsics.h new file mode 100644 index 0000000000..b608ad6819 --- /dev/null +++ b/include/hip/hcc_detail/llvm_intrinsics.h @@ -0,0 +1,46 @@ +/* +Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved. + +Permission is hereby granted, free of charge, to any person obtaining a copy +of this software and associated documentation files (the "Software"), to deal +in the Software without restriction, including without limitation the rights +to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +copies of the Software, and to permit persons to whom the Software is +furnished to do so, subject to the following conditions: + +The above copyright notice and this permission notice shall be included in +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, +OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN +THE SOFTWARE. +*/ + +/** + * @file hcc_detail/llvm_intrinsics.h + * @brief Contains declarations for wrapper functions for llvm intrinsics + * like llvm.amdgcn.s.barrier. + */ + +#ifndef HIP_INCLUDE_HIP_HCC_DETAIL_LLVM_INTRINSICS_H +#define HIP_INCLUDE_HIP_HCC_DETAIL_LLVM_INTRINSICS_H + +#include "hip/hcc_detail/host_defines.h" + +__device__ +unsigned long __llvm_amdgcn_s_memrealtime(void) __asm("llvm.amdgcn.s.memrealtime"); + +__device__ +unsigned __llvm_amdgcn_s_getreg(unsigned) __asm("llvm.amdgcn.s.getreg"); + +__device__ +unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); + +__device__ +void __llvm_amdgcn_s_barrier() __asm("llvm.amdgcn.s.barrier"); + +#endif