From c453b42bff45800ab04e10231dcf63a6acddc876 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 12 Jun 2018 22:05:59 +0000 Subject: [PATCH 01/21] Add hipclang amdgcn functions These are moving from hipclang in device library to hip headers. These are required for the functionality of HIPclang project. --- include/hip/hcc_detail/hip_runtime.h | 197 ++++++++++++++++++++++++++- src/device_util.cpp | 9 +- 2 files changed, 192 insertions(+), 14 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index c62c85df64..34650d728a 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -189,8 +189,16 @@ extern int HIP_TRACE_API; static constexpr int warpSize = 64; #define clock_t long long int -__device__ long long int clock64(); -__device__ clock_t clock(); +__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(); } + +__device__ +inline +clock_t __clock() { return (clock_t)__llvm_amdgcn_s_memrealtime(); } // abort __device__ void abort(); @@ -200,6 +208,45 @@ __device__ int __all(int input); __device__ int __any(int input); __device__ unsigned long long int __ballot(int input); +__device__ +inline +int64_t __ballot64(int a) { + int64_t s; + // define i64 @__ballot64(i32 %a) #0 { + // %b = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %a) #1 + // ret i64 %b + // } + __asm("v_cmp_ne_i32_e64 $0, 0, $1" : "=s"(s) : "v"(a)); + return s; +} + +// hip.amdgcn.bc - lanemask +extern "C" __device__ int32_t __ockl_activelane_u32(void); + +__device__ +inline +int64_t __lanemask_gt() +{ + int32_t activelane = __ockl_activelane_u32(); + int64_t ballot = __ballot64(1); + if (activelane != 63) { + int64_t tmp = (~0UL) << (activelane + 1); + return tmp & ballot; + } + return 0; +} + +__device__ +inline +int64_t __lanemask_lt() +{ + int32_t activelane = __ockl_activelane_u32(); + int64_t ballot = __ballot64(1); + if (activelane == 0) + return 0; + return ballot; +} + #if __HIP_ARCH_GFX701__ == 0 // warp shuffle functions @@ -238,8 +285,11 @@ __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); -__device__ void* __get_dynamicgroupbaseptr(); +extern "C" __device__ void* get_dynamic_group_segment_base_pointer(); +__device__ +inline +void* __get_dynamicgroupbaseptr() { return get_dynamic_group_segment_base_pointer(); } /** * CUDA 8 device function features @@ -315,6 +365,11 @@ __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__ #if defined __HCC__ @@ -572,7 +627,9 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; __DEVICE__ void __device_trap() __asm("llvm.trap"); -__DEVICE__ void inline __assert_fail(const char * __assertion, +__DEVICE__ +inline +void __assert_fail(const char * __assertion, const char *__file, unsigned int __line, const char *__function) @@ -581,8 +638,136 @@ __DEVICE__ void inline __assert_fail(const char * __assertion, __device_trap(); } -extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); -extern "C" __device__ void *__amdgcn_get_dynamicgroupbaseptr(); +__DEVICE__ +inline +void __assertfail(const char * __assertion, + const char *__file, + unsigned int __line, + const char *__function, + size_t charsize) +{ + // ignore all the args for now. + __device_trap(); +} + +// hip.amdgcn.bc - sync threads +// extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); +#define CLK_LOCAL_MEM_FENCE 0x01 +#define local __attribute__((address_space(3))) + +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; + +extern "C" __device__ __attribute__((overloadable)) +void atomic_work_item_fence(cl_mem_fence_flags, memory_order, memory_scope); + +__device__ +inline +static void hc_work_group_barrier(cl_mem_fence_flags flags, memory_scope scope) +{ + if (flags) { + atomic_work_item_fence(flags, memory_order_release, scope); + __builtin_amdgcn_s_barrier(); + atomic_work_item_fence(flags, memory_order_acquire, scope); + } else { + __builtin_amdgcn_s_barrier(); + } +} + +__device__ +inline +static void hc_barrier(int n) +{ + hc_work_group_barrier((cl_mem_fence_flags)n, memory_scope_work_group); +} + +__device__ +inline +__attribute__((noduplicate)) +void __syncthreads() +{ + hc_barrier(CLK_LOCAL_MEM_FENCE); +} + + +__device__ unsigned __llvm_amdgcn_s_getreg(unsigned) __asm("llvm.amdgcn.s.getreg"); + +__device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); + +__device__ inline static local char* __to_local(unsigned x) { return (local char*)x; } + +__device__ inline void *__amdgcn_get_dynamicgroupbaseptr() { +#if 0 + // Get group segment base pointer. + char* base = __llvm_amdgcn_s_getreg(14342) << 8); + base += __llvm_amdgcn_groupstaticsize(); + return base; +#endif + return __get_dynamicgroupbaseptr(); +} + +// hip.amdgcn.bc - device routine +/* + HW_ID Register bit structure + WAVE_ID 3:0 Wave buffer slot number. 0-9. + SIMD_ID 5:4 SIMD which the wave is assigned to within the CU. + PIPE_ID 7:6 Pipeline from which the wave was dispatched. + CU_ID 11:8 Compute Unit the wave is assigned to. + SH_ID 12 Shader Array (within an SE) the wave is assigned to. + SE_ID 14:13 Shader Engine the wave is assigned to. + TG_ID 19:16 Thread-group ID + VM_ID 23:20 Virtual Memory ID + QUEUE_ID 26:24 Queue from which this wave was dispatched. + STATE_ID 29:27 State ID (graphics only, not compute). + ME_ID 31:30 Micro-engine ID. + */ + +#define HW_ID 4 + +#define HW_ID_CU_ID_SIZE 4 +#define HW_ID_CU_ID_OFFSET 8 + +#define HW_ID_SE_ID_SIZE 2 +#define HW_ID_SE_ID_OFFSET 13 + +/* + Encoding of parameter bitmask + HW_ID 5:0 HW_ID + OFFSET 10:6 Range: 0..31 + SIZE 15:11 Range: 1..32 + */ + +#define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG + +__device__ +inline +unsigned __smid(void) +{ + unsigned cu_id = __llvm_amdgcn_s_getreg( + GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID)); + unsigned se_id = __llvm_amdgcn_s_getreg( + GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID)); + + /* Each shader engine has 16 CU */ + return (se_id << HW_ID_CU_ID_SIZE) + cu_id; +} // Macro to replace extern __shared__ declarations // to local variable definitions diff --git a/src/device_util.cpp b/src/device_util.cpp index 613e35f0cc..a3386ba14d 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -144,9 +144,6 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) { return dst; } -__device__ long long int clock64() { return (long long int)hc::__cycle_u64(); }; -__device__ clock_t clock() { return (clock_t)hc::__cycle_u64(); }; - // abort __device__ void abort() { return hc::abort(); } @@ -203,11 +200,7 @@ __host__ __device__ int max(int arg1, int arg2) { return (int)(hc::precise_math::fmax((float)arg1, (float)arg2)); } -__device__ void* __get_dynamicgroupbaseptr() { - return hc::get_dynamic_group_segment_base_pointer(); -} - __host__ void* __get_dynamicgroupbaseptr() { return nullptr; } -__device__ void __threadfence_system(void) { std::atomic_thread_fence(std::memory_order_seq_cst); } \ No newline at end of file +__device__ void __threadfence_system(void) { std::atomic_thread_fence(std::memory_order_seq_cst); } From cfa8fc1ca554d05c9538cc481523cc70954c1d4e Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 13 Jun 2018 15:59:45 +0000 Subject: [PATCH 02/21] Add __llvm_fence funcs and __ prefixes --- include/hip/hcc_detail/hip_runtime.h | 52 ++++++++++++++++++++-------- 1 file changed, 37 insertions(+), 15 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 34650d728a..b9c7ea603c 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -652,40 +652,62 @@ void __assertfail(const char * __assertion, // hip.amdgcn.bc - sync threads // extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); -#define CLK_LOCAL_MEM_FENCE 0x01 -#define local __attribute__((address_space(3))) +#define __CLK_LOCAL_MEM_FENCE 0x01 +#define __local __attribute__((address_space(3))) -typedef unsigned cl_mem_fence_flags; +typedef unsigned __cl_mem_fence_flags; -typedef enum memory_scope { +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; +} __memory_scope; // enum values aligned with what clang uses in EmitAtomicExpr() -typedef enum memory_order +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; +} __memory_order; -extern "C" __device__ __attribute__((overloadable)) -void atomic_work_item_fence(cl_mem_fence_flags, memory_order, memory_scope); +// __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 hc_work_group_barrier(cl_mem_fence_flags flags, memory_scope scope) +static void hc_work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) { if (flags) { - atomic_work_item_fence(flags, memory_order_release, scope); + switch(scope) { + case memory_scope_work_item: break; + case memory_scope_sub_group: __llvm_fence_rel_sg(); break; + case memory_scope_work_group: __llvm_fence_rel_wg(); break; + case memory_scope_device: __llvm_fence_rel_dev(); break; + case memory_scope_all_svm_devices: __llvm_fence_rel_sys(); break; + } + //atomic_work_item_fence(flags, memory_order_release, scope); __builtin_amdgcn_s_barrier(); - atomic_work_item_fence(flags, memory_order_acquire, scope); + //atomic_work_item_fence(flags, memory_order_acquire, scope); + switch(scope) { + case memory_scope_work_item: break; + case memory_scope_sub_group: __llvm_fence_acq_sg(); break; + case memory_scope_work_group: __llvm_fence_acq_wg(); break; + case memory_scope_device: __llvm_fence_acq_dev(); break; + case memory_scope_all_svm_devices: __llvm_fence_acq_sys(); break; + } } else { __builtin_amdgcn_s_barrier(); } @@ -695,7 +717,7 @@ __device__ inline static void hc_barrier(int n) { - hc_work_group_barrier((cl_mem_fence_flags)n, memory_scope_work_group); + hc_work_group_barrier((__cl_mem_fence_flags)n, memory_scope_work_group); } __device__ @@ -703,7 +725,7 @@ inline __attribute__((noduplicate)) void __syncthreads() { - hc_barrier(CLK_LOCAL_MEM_FENCE); + hc_barrier(__CLK_LOCAL_MEM_FENCE); } @@ -711,7 +733,7 @@ __device__ unsigned __llvm_amdgcn_s_getreg(unsigned) __asm("llvm.amdgcn.s.getreg __device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); -__device__ inline static local char* __to_local(unsigned x) { return (local char*)x; } +__device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } __device__ inline void *__amdgcn_get_dynamicgroupbaseptr() { #if 0 From 871cfc6fd223fe3eeff18e02df8e4f2b70360bda Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 13 Jun 2018 16:06:58 +0000 Subject: [PATCH 03/21] Add prefix __ to memory scope and order --- include/hip/hcc_detail/hip_runtime.h | 42 ++++++++++++++-------------- 1 file changed, 21 insertions(+), 21 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index b9c7ea603c..58bbc1a91d 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -658,21 +658,21 @@ void __assertfail(const char * __assertion, 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_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_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; // __llvm_fence* functions from device-libs/irif/src/fence.ll @@ -692,21 +692,21 @@ static void hc_work_group_barrier(__cl_mem_fence_flags flags, __memory_scope sco { if (flags) { switch(scope) { - case memory_scope_work_item: break; - case memory_scope_sub_group: __llvm_fence_rel_sg(); break; - case memory_scope_work_group: __llvm_fence_rel_wg(); break; - case memory_scope_device: __llvm_fence_rel_dev(); break; - case memory_scope_all_svm_devices: __llvm_fence_rel_sys(); break; + case __memory_scope_work_item: break; + case __memory_scope_sub_group: __llvm_fence_rel_sg(); break; + case __memory_scope_work_group: __llvm_fence_rel_wg(); break; + case __memory_scope_device: __llvm_fence_rel_dev(); break; + case __memory_scope_all_svm_devices: __llvm_fence_rel_sys(); break; } //atomic_work_item_fence(flags, memory_order_release, scope); __builtin_amdgcn_s_barrier(); //atomic_work_item_fence(flags, memory_order_acquire, scope); switch(scope) { - case memory_scope_work_item: break; - case memory_scope_sub_group: __llvm_fence_acq_sg(); break; - case memory_scope_work_group: __llvm_fence_acq_wg(); break; - case memory_scope_device: __llvm_fence_acq_dev(); break; - case memory_scope_all_svm_devices: __llvm_fence_acq_sys(); break; + case __memory_scope_work_item: break; + case __memory_scope_sub_group: __llvm_fence_acq_sg(); break; + case __memory_scope_work_group: __llvm_fence_acq_wg(); break; + case __memory_scope_device: __llvm_fence_acq_dev(); break; + case __memory_scope_all_svm_devices: __llvm_fence_acq_sys(); break; } } else { __builtin_amdgcn_s_barrier(); @@ -717,7 +717,7 @@ __device__ inline static void hc_barrier(int n) { - hc_work_group_barrier((__cl_mem_fence_flags)n, memory_scope_work_group); + hc_work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group); } __device__ From 27f600b425becf76009550ec89613480be4381f6 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 13 Jun 2018 20:40:10 +0000 Subject: [PATCH 04/21] Add get_dynamicgroupbaseptr def and remove hc_ --- include/hip/hcc_detail/hip_runtime.h | 49 ++++++++++++++-------------- 1 file changed, 24 insertions(+), 25 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 58bbc1a91d..1589f19395 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -285,11 +285,29 @@ __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); -extern "C" __device__ void* get_dynamic_group_segment_base_pointer(); +// 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() { return get_dynamic_group_segment_base_pointer(); } +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); +} + +__device__ inline void *__amdgcn_get_dynamicgroupbaseptr() { + return __get_dynamicgroupbaseptr(); +} /** * CUDA 8 device function features @@ -653,8 +671,6 @@ void __assertfail(const char * __assertion, // hip.amdgcn.bc - sync threads // extern "C" __device__ __attribute__((noduplicate)) void __syncthreads(); #define __CLK_LOCAL_MEM_FENCE 0x01 -#define __local __attribute__((address_space(3))) - typedef unsigned __cl_mem_fence_flags; typedef enum __memory_scope { @@ -688,7 +704,7 @@ extern "C" __device__ void __llvm_fence_rel_sys(void); __device__ inline -static void hc_work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) +static void __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) { if (flags) { switch(scope) { @@ -715,9 +731,9 @@ static void hc_work_group_barrier(__cl_mem_fence_flags flags, __memory_scope sco __device__ inline -static void hc_barrier(int n) +static void __barrier(int n) { - hc_work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group); + __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group); } __device__ @@ -725,24 +741,7 @@ inline __attribute__((noduplicate)) void __syncthreads() { - hc_barrier(__CLK_LOCAL_MEM_FENCE); -} - - -__device__ unsigned __llvm_amdgcn_s_getreg(unsigned) __asm("llvm.amdgcn.s.getreg"); - -__device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); - -__device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } - -__device__ inline void *__amdgcn_get_dynamicgroupbaseptr() { -#if 0 - // Get group segment base pointer. - char* base = __llvm_amdgcn_s_getreg(14342) << 8); - base += __llvm_amdgcn_groupstaticsize(); - return base; -#endif - return __get_dynamicgroupbaseptr(); + __barrier(__CLK_LOCAL_MEM_FENCE); } // hip.amdgcn.bc - device routine From cfe37484c9e35389a6a490b8f4efe7d5f1735690 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 14 Jun 2018 17:49:35 +0000 Subject: [PATCH 05/21] 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 From fe4e6c53fce33b71bcbe10e3a4d0f157a212cb9d Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 15 Jun 2018 23:06:40 +0000 Subject: [PATCH 06/21] Move hipclang funcs into corresponding headers --- include/hip/hcc_detail/device_functions.h | 256 ++++++++++++++++++++++ include/hip/hcc_detail/hip_runtime.h | 229 ------------------- include/hip/hcc_detail/llvm_intrinsics.h | 9 - 3 files changed, 256 insertions(+), 238 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index ae1b96c979..8bae5325fd 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -24,8 +24,12 @@ THE SOFTWARE. #define HIP_INCLUDE_HIP_HCC_DETAIL_DEVICE_FUNCTIONS_H #include "host_defines.h" +#include "math_fwd.h" +#include #include +#include +#include extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int); extern "C" __device__ signed int __hip_hc_ir_mul24_int(signed int, signed int); @@ -209,5 +213,257 @@ __device__ char4 __hip_hc_add8pk(char4, char4); __device__ char4 __hip_hc_sub8pk(char4, char4); __device__ char4 __hip_hc_mul8pk(char4, char4); +#if defined(__HCC__) +#define __HCC_OR_HIP_CLANG__ 1 +#elif defined(__clang__) && defined(__HIP__) +#define __HCC_OR_HIP_CLANG__ 1 +#else +#define __HCC_OR_HIP_CLANG__ 0 +#endif + +#ifdef __HCC_OR_HIP_CLANG__ + +#ifdef __HIP_DEVICE_COMPILE__ + +// Clock functions +__device__ +inline +long long int __clock64() { return (long long int) __builtin_amdgcn_s_memrealtime(); } + +__device__ +inline +long long int __clock() { return (long long int) __builtin_amdgcn_s_memrealtime(); } + +// hip.amdgcn.bc - named sync +__device__ +inline +void __named_sync(int a, int b) { __builtin_amdgcn_s_barrier(); } + +#endif // __HIP_DEVICE_COMPILE__ + +// warp vote function __all __any __ballot +__device__ +int __all(int input); +__device__ +int __any(int input); +__device__ +unsigned long long int __ballot(int input); + +__device__ +inline +uint64_t __ballot64(int a) { + int64_t s; + // define i64 @__ballot64(i32 %a) #0 { + // %b = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %a) #1 + // ret i64 %b + // } + __asm("v_cmp_ne_i32_e64 $0, 0, $1" : "=s"(s) : "v"(a)); + return s; +} + +// hip.amdgcn.bc - lanemask +__device__ +inline +int64_t __lanemask_gt() +{ + int32_t activelane = __ockl_activelane_u32(); + int64_t ballot = __ballot64(1); + if (activelane != 63) { + int64_t tmp = (~0UL) << (activelane + 1); + return tmp & ballot; + } + return 0; +} + +__device__ +inline +int64_t __lanemask_lt() +{ + int32_t activelane = __ockl_activelane_u32(); + int64_t ballot = __ballot64(1); + if (activelane == 0) + return 0; + return ballot; +} + +__device__ +inline +void* __get_dynamicgroupbaseptr() +{ + // Get group segment base pointer. + return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); +} + +__device__ +inline +void *__amdgcn_get_dynamicgroupbaseptr() { + return __get_dynamicgroupbaseptr(); +} + +#endif // __HCC_OR_HIP_CLANG__ + +#ifdef __HCC__ + +/** + * extern __shared__ + */ + +// Macro to replace extern __shared__ declarations +// to local variable definitions +#define HIP_DYNAMIC_SHARED(type, var) type* var = (type*)__get_dynamicgroupbaseptr(); + +#define HIP_DYNAMIC_SHARED_ATTRIBUTE + + +#elif defined(__clang__) && defined(__HIP__) + +#pragma push_macro("__DEVICE__") +#define __DEVICE__ extern "C" __device__ __attribute__((always_inline)) \ + __attribute__((weak)) + +__DEVICE__ +inline +void __assert_fail(const char * __assertion, + const char *__file, + unsigned int __line, + const char *__function) +{ + // Ignore all the args for now. + __builtin_trap(); +} + +__DEVICE__ +inline +void __assertfail(const char * __assertion, + const char *__file, + unsigned int __line, + const char *__function, + size_t charsize) +{ + // ignore all the args for now. + __builtin_trap(); +} + +// 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 __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) +{ + if (flags) { + switch(scope) { + case __memory_scope_work_item: break; + case __memory_scope_sub_group: __llvm_fence_rel_sg(); break; + case __memory_scope_work_group: __llvm_fence_rel_wg(); break; + case __memory_scope_device: __llvm_fence_rel_dev(); break; + case __memory_scope_all_svm_devices: __llvm_fence_rel_sys(); break; + } + //atomic_work_item_fence(flags, memory_order_release, scope); + __builtin_amdgcn_s_barrier(); + //atomic_work_item_fence(flags, memory_order_acquire, scope); + switch(scope) { + case __memory_scope_work_item: break; + case __memory_scope_sub_group: __llvm_fence_acq_sg(); break; + case __memory_scope_work_group: __llvm_fence_acq_wg(); break; + case __memory_scope_device: __llvm_fence_acq_dev(); break; + case __memory_scope_all_svm_devices: __llvm_fence_acq_sys(); break; + } + } else { + __builtin_amdgcn_s_barrier(); + } +} + +__device__ +inline +static void __barrier(int n) +{ + __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group); +} + +__device__ +inline +__attribute__((noduplicate)) +void __syncthreads() +{ + __barrier(__CLK_LOCAL_MEM_FENCE); +} + +// hip.amdgcn.bc - device routine +/* + HW_ID Register bit structure + WAVE_ID 3:0 Wave buffer slot number. 0-9. + SIMD_ID 5:4 SIMD which the wave is assigned to within the CU. + PIPE_ID 7:6 Pipeline from which the wave was dispatched. + CU_ID 11:8 Compute Unit the wave is assigned to. + SH_ID 12 Shader Array (within an SE) the wave is assigned to. + SE_ID 14:13 Shader Engine the wave is assigned to. + TG_ID 19:16 Thread-group ID + VM_ID 23:20 Virtual Memory ID + QUEUE_ID 26:24 Queue from which this wave was dispatched. + STATE_ID 29:27 State ID (graphics only, not compute). + ME_ID 31:30 Micro-engine ID. + */ + +#define HW_ID 4 + +#define HW_ID_CU_ID_SIZE 4 +#define HW_ID_CU_ID_OFFSET 8 + +#define HW_ID_SE_ID_SIZE 2 +#define HW_ID_SE_ID_OFFSET 13 + +/* + Encoding of parameter bitmask + HW_ID 5:0 HW_ID + OFFSET 10:6 Range: 0..31 + SIZE 15:11 Range: 1..32 + */ + +#define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG + +__device__ +inline +unsigned __smid(void) +{ + unsigned cu_id = __builtin_amdgcn_s_getreg( + GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID)); + unsigned se_id = __builtin_amdgcn_s_getreg( + GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID)); + + /* Each shader engine has 16 CU */ + return (se_id << HW_ID_CU_ID_SIZE) + cu_id; +} + +#pragma push_macro("__DEVICE__") + +// Macro to replace extern __shared__ declarations +// to local variable definitions +#define HIP_DYNAMIC_SHARED(type, var) \ + type* var = (type*)__amdgcn_get_dynamicgroupbaseptr(); + +#define HIP_DYNAMIC_SHARED_ATTRIBUTE + + +#endif //defined(__clang__) && defined(__HIP__) #endif diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 4cd41a0c86..18b04daf77 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -62,8 +62,6 @@ THE SOFTWARE. #define CUDA_SUCCESS hipSuccess #include -#include -#include #endif // __HCC_OR_HIP_CLANG__ #if __HCC__ @@ -190,60 +188,9 @@ extern int HIP_TRACE_API; //#define warpSize hc::__wavesize() static constexpr int warpSize = 64; -#define clock_t long long int -__device__ -inline -long long int __clock64() { return (long long int)__llvm_amdgcn_s_memrealtime(); } - -__device__ -inline -clock_t __clock() { return (clock_t)__llvm_amdgcn_s_memrealtime(); } - // abort __device__ void abort(); -// warp vote function __all __any __ballot -__device__ int __all(int input); -__device__ int __any(int input); -__device__ unsigned long long int __ballot(int input); - -__device__ -inline -int64_t __ballot64(int a) { - int64_t s; - // define i64 @__ballot64(i32 %a) #0 { - // %b = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %a) #1 - // ret i64 %b - // } - __asm("v_cmp_ne_i32_e64 $0, 0, $1" : "=s"(s) : "v"(a)); - return s; -} - -// hip.amdgcn.bc - lanemask -__device__ -inline -int64_t __lanemask_gt() -{ - int32_t activelane = __ockl_activelane_u32(); - int64_t ballot = __ballot64(1); - if (activelane != 63) { - int64_t tmp = (~0UL) << (activelane + 1); - return tmp & ballot; - } - return 0; -} - -__device__ -inline -int64_t __lanemask_lt() -{ - int32_t activelane = __ockl_activelane_u32(); - int64_t ballot = __ballot64(1); - if (activelane == 0) - return 0; - return ballot; -} - #if __HIP_ARCH_GFX701__ == 0 // warp shuffle functions @@ -283,20 +230,6 @@ __host__ __device__ int min(int arg1, int arg2); __host__ __device__ int max(int arg1, int arg2); -__device__ -inline -void* __get_dynamicgroupbaseptr() -{ - // Get group segment base pointer. - return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); -} - -__device__ -inline -void *__amdgcn_get_dynamicgroupbaseptr() { - return __get_dynamicgroupbaseptr(); -} - /** * CUDA 8 device function features @@ -371,9 +304,6 @@ __device__ void __threadfence_system(void); * @} */ -// hip.amdgcn.bc - named sync -__device__ inline void __named_sync(int a, int b) { __llvm_amdgcn_s_barrier(); } - #endif // __HCC_OR_HIP_CLANG__ #if defined __HCC__ @@ -496,17 +426,6 @@ extern void ihipPostLaunchKernel(const char* kernelName, hipStream_t stream, gri #endif //__HCC_CPP__ -/** - * extern __shared__ - */ - -// Macro to replace extern __shared__ declarations -// to local variable definitions -#define HIP_DYNAMIC_SHARED(type, var) type* var = (type*)__get_dynamicgroupbaseptr(); - -#define HIP_DYNAMIC_SHARED_ATTRIBUTE - - /** * @defgroup HIP-ENV HIP Environment Variables * @{ @@ -625,154 +544,6 @@ extern const __device__ __attribute__((weak)) __hip_builtin_gridDim_t gridDim; #define hipGridDim_y gridDim.y #define hipGridDim_z gridDim.z -#pragma push_macro("__DEVICE__") -#define __DEVICE__ extern "C" __device__ __attribute__((always_inline)) \ - __attribute__((weak)) - -__DEVICE__ void __device_trap() __asm("llvm.trap"); - -__DEVICE__ -inline -void __assert_fail(const char * __assertion, - const char *__file, - unsigned int __line, - const char *__function) -{ - // Ignore all the args for now. - __device_trap(); -} - -__DEVICE__ -inline -void __assertfail(const char * __assertion, - const char *__file, - unsigned int __line, - const char *__function, - size_t charsize) -{ - // ignore all the args for now. - __device_trap(); -} - -// 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 __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) -{ - if (flags) { - switch(scope) { - case __memory_scope_work_item: break; - case __memory_scope_sub_group: __llvm_fence_rel_sg(); break; - case __memory_scope_work_group: __llvm_fence_rel_wg(); break; - case __memory_scope_device: __llvm_fence_rel_dev(); break; - case __memory_scope_all_svm_devices: __llvm_fence_rel_sys(); break; - } - //atomic_work_item_fence(flags, memory_order_release, scope); - __builtin_amdgcn_s_barrier(); - //atomic_work_item_fence(flags, memory_order_acquire, scope); - switch(scope) { - case __memory_scope_work_item: break; - case __memory_scope_sub_group: __llvm_fence_acq_sg(); break; - case __memory_scope_work_group: __llvm_fence_acq_wg(); break; - case __memory_scope_device: __llvm_fence_acq_dev(); break; - case __memory_scope_all_svm_devices: __llvm_fence_acq_sys(); break; - } - } else { - __builtin_amdgcn_s_barrier(); - } -} - -__device__ -inline -static void __barrier(int n) -{ - __work_group_barrier((__cl_mem_fence_flags)n, __memory_scope_work_group); -} - -__device__ -inline -__attribute__((noduplicate)) -void __syncthreads() -{ - __barrier(__CLK_LOCAL_MEM_FENCE); -} - -// hip.amdgcn.bc - device routine -/* - HW_ID Register bit structure - WAVE_ID 3:0 Wave buffer slot number. 0-9. - SIMD_ID 5:4 SIMD which the wave is assigned to within the CU. - PIPE_ID 7:6 Pipeline from which the wave was dispatched. - CU_ID 11:8 Compute Unit the wave is assigned to. - SH_ID 12 Shader Array (within an SE) the wave is assigned to. - SE_ID 14:13 Shader Engine the wave is assigned to. - TG_ID 19:16 Thread-group ID - VM_ID 23:20 Virtual Memory ID - QUEUE_ID 26:24 Queue from which this wave was dispatched. - STATE_ID 29:27 State ID (graphics only, not compute). - ME_ID 31:30 Micro-engine ID. - */ - -#define HW_ID 4 - -#define HW_ID_CU_ID_SIZE 4 -#define HW_ID_CU_ID_OFFSET 8 - -#define HW_ID_SE_ID_SIZE 2 -#define HW_ID_SE_ID_OFFSET 13 - -/* - Encoding of parameter bitmask - HW_ID 5:0 HW_ID - OFFSET 10:6 Range: 0..31 - SIZE 15:11 Range: 1..32 - */ - -#define GETREG_IMMED(SZ,OFF,REG) (SZ << 11) | (OFF << 6) | REG - -__device__ -inline -unsigned __smid(void) -{ - unsigned cu_id = __llvm_amdgcn_s_getreg( - GETREG_IMMED(HW_ID_CU_ID_SIZE, HW_ID_CU_ID_OFFSET, HW_ID)); - unsigned se_id = __llvm_amdgcn_s_getreg( - GETREG_IMMED(HW_ID_SE_ID_SIZE, HW_ID_SE_ID_OFFSET, HW_ID)); - - /* Each shader engine has 16 CU */ - return (se_id << HW_ID_CU_ID_SIZE) + cu_id; -} - -// Macro to replace extern __shared__ declarations -// to local variable definitions -#define HIP_DYNAMIC_SHARED(type, var) \ - type* var = (type*)__amdgcn_get_dynamicgroupbaseptr(); - -#define HIP_DYNAMIC_SHARED_ATTRIBUTE - -#pragma push_macro("__DEVICE__") - #include #endif diff --git a/include/hip/hcc_detail/llvm_intrinsics.h b/include/hip/hcc_detail/llvm_intrinsics.h index b608ad6819..2c7819b535 100644 --- a/include/hip/hcc_detail/llvm_intrinsics.h +++ b/include/hip/hcc_detail/llvm_intrinsics.h @@ -31,16 +31,7 @@ THE SOFTWARE. #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 From e02fc7e680059047ab6bd7db29c39de017e0cf84 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 19 Jun 2018 19:22:31 +0000 Subject: [PATCH 07/21] Implement device_functions.cpp into HIP headers Move all Integer Intrinsics, device_functions.cpp definitions and HIP specific device functions into HIP headers. Implement the device functions using llvm_intrinsics and device-libs functions instead of calling hc::__* functions. Remove device_functions.cpp since everything is now defined in header. --- CMakeLists.txt | 3 +- include/hip/hcc_detail/device_functions.h | 475 ++++++++++++++---- include/hip/hcc_detail/device_library_decls.h | 5 + include/hip/hcc_detail/llvm_intrinsics.h | 29 ++ src/device_functions.cpp | 387 -------------- 5 files changed, 422 insertions(+), 477 deletions(-) delete mode 100644 src/device_functions.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 537764548f..b649b20c21 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -194,8 +194,7 @@ if(HIP_PLATFORM STREQUAL "hcc") set(SOURCE_FILES_DEVICE src/device_util.cpp - src/hip_ldg.cpp - src/device_functions.cpp) + src/hip_ldg.cpp) execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 8bae5325fd..b147cd9b80 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -31,35 +31,142 @@ THE SOFTWARE. #include #include +typedef unsigned long ulong; +typedef unsigned int uint; + extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int); extern "C" __device__ signed int __hip_hc_ir_mul24_int(signed int, signed int); extern "C" __device__ signed int __hip_hc_ir_mulhi_int(signed int, signed int); extern "C" __device__ unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int); extern "C" __device__ unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int); +/* +Integer Intrinsics +*/ + // integer intrinsic function __poc __clz __ffs __brev -__device__ unsigned int __brev(unsigned int x); -__device__ unsigned long long int __brevll(unsigned long long int x); -__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s); -__device__ unsigned int __clz(int x); -__device__ unsigned int __clzll(long long int x); -__device__ unsigned int __ffs(int x); -__device__ unsigned int __ffsll(long long int x); +__device__ static inline unsigned int __popc(unsigned int input) { + return __builtin_popcount(input); +} +__device__ static inline unsigned int __popcll(unsigned long long int input) { + return __builtin_popcountl(input); +} + +__device__ static inline unsigned int __clz(unsigned int input) { +#ifdef NVCC_COMPAT + return input == 0 ? 32 : __builtin_clz(input); +#else + return input == 0 ? -1 : __builtin_clz(input); +#endif +} + +__device__ static inline unsigned int __clzll(unsigned long long int input) { +#ifdef NVCC_COMPAT + return input == 0 ? 64 : ( input == 0 ? -1 : __builtin_clzl(input) ); +#else + return input == 0 ? -1 : __builtin_clzl(input); +#endif +} + +__device__ static inline unsigned int __clz(int input) { +#ifdef NVCC_COMPAT + return input == 0 ? 32 : ( input > 0 ? __builtin_clz(input) : __builtin_clz(~input) ); +#else + if (input == 0) return -1; + return input > 0 ? __builtin_clz(input) : __builtin_clz(~input); +#endif +} + +__device__ static inline unsigned int __clzll(long long int input) { +#ifdef NVCC_COMPAT + return input == 0 ? 64 : input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input); +#else + if (input == 0) return -1; + return input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input); +#endif +} + +__device__ static inline unsigned int __ffs(unsigned int input) { +#ifdef NVCC_COMPAT + return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1; +#else + return input == 0 ? -1 : __builtin_ctz(input); +#endif +} + +__device__ static inline unsigned int __ffsll(unsigned long long int input) { +#ifdef NVCC_COMPAT + return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; +#else + return input == 0 ? -1 : __builtin_ctzl(input); +#endif +} + +__device__ static inline unsigned int __ffs(int input) { +#ifdef NVCC_COMPAT + return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1; +#else + return input == 0 ? -1 : __builtin_ctz(input); +#endif +} + +__device__ static inline unsigned int __ffsll(long long int input) { +#ifdef NVCC_COMPAT + return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1; +#else + return input == 0 ? -1 : __builtin_ctzl(input); +#endif +} + +__device__ static inline unsigned int __brev(unsigned int input) { return __llvm_bitrev_b32(input); } + +__device__ static inline unsigned long long int __brevll(unsigned long long int input) { + return __llvm_bitrev_b64(input); +} + +__device__ static unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s); __device__ static unsigned int __hadd(int x, int y); __device__ static int __mul24(int x, int y); -__device__ long long int __mul64hi(long long int x, long long int y); +__device__ static long long int __mul64hi(long long int x, long long int y); __device__ static int __mulhi(int x, int y); -__device__ unsigned int __popc(unsigned int x); -__device__ unsigned int __popcll(unsigned long long int x); __device__ static int __rhadd(int x, int y); __device__ static unsigned int __sad(int x, int y, int z); __device__ static unsigned int __uhadd(unsigned int x, unsigned int y); __device__ static int __umul24(unsigned int x, unsigned int y); -__device__ unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y); +__device__ static unsigned long long int __umul64hi(unsigned long long int x, unsigned long long int y); __device__ static unsigned int __umulhi(unsigned int x, unsigned int y); __device__ static unsigned int __urhadd(unsigned int x, unsigned int y); __device__ static unsigned int __usad(unsigned int x, unsigned int y, unsigned int z); +struct ucharHolder { + union { + unsigned char c[4]; + unsigned int ui; + }; +} __attribute__((aligned(4))); + +struct uchar2Holder { + union { + unsigned int ui[2]; + unsigned char c[8]; + }; +} __attribute__((aligned(8))); + +__device__ +static inline unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) { + struct uchar2Holder cHoldVal; + struct ucharHolder cHoldKey; + struct ucharHolder cHoldOut; + cHoldKey.ui = s; + cHoldVal.ui[0] = x; + cHoldVal.ui[1] = y; + cHoldOut.c[0] = cHoldVal.c[cHoldKey.c[0]]; + cHoldOut.c[1] = cHoldVal.c[cHoldKey.c[1]]; + cHoldOut.c[2] = cHoldVal.c[cHoldKey.c[2]]; + cHoldOut.c[3] = cHoldVal.c[cHoldKey.c[3]]; + return cHoldOut.ui; +} + __device__ static inline unsigned int __hadd(int x, int y) { int z = x + y; int sign = z & 0x8000000; @@ -67,6 +174,20 @@ __device__ static inline unsigned int __hadd(int x, int y) { return ((value) >> 1 || sign); } __device__ static inline int __mul24(int x, int y) { return __hip_hc_ir_mul24_int(x, y); } + +__device__ static inline long long __mul64hi(long long int x, long long int y) { + ulong x0 = (ulong)x & 0xffffffffUL; + long x1 = x >> 32; + ulong y0 = (ulong)y & 0xffffffffUL; + long y1 = y >> 32; + ulong z0 = x0*y0; + long t = x1*y0 + (z0 >> 32); + long z1 = t & 0xffffffffL; + long z2 = t >> 32; + z1 = x0*y1 + z1; + return x1*y1 + z2 + (z1 >> 32); +} + __device__ static inline int __mulhi(int x, int y) { return __hip_hc_ir_mulhi_int(x, y); } __device__ static inline int __rhadd(int x, int y) { int z = x + y + 1; @@ -83,6 +204,21 @@ __device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) { __device__ static inline int __umul24(unsigned int x, unsigned int y) { return __hip_hc_ir_umul24_int(x, y); } + +__device__ +static inline unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) { + ulong x0 = x & 0xffffffffUL; + ulong x1 = x >> 32; + ulong y0 = y & 0xffffffffUL; + ulong y1 = y >> 32; + ulong z0 = x0*y0; + ulong t = x1*y0 + (z0 >> 32); + ulong z1 = t & 0xffffffffUL; + ulong z2 = t >> 32; + z1 = x0*y1 + z1; + return x1*y1 + z2 + (z1 >> 32); +} + __device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y) { return __hip_hc_ir_umulhi_int(x, y); } @@ -93,45 +229,158 @@ __device__ static inline unsigned int __usad(unsigned int x, unsigned int y, uns return __hip_hc_ir_usad_int(x, y, z); } -extern __device__ __attribute__((const)) unsigned int __mbcnt_lo(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.lo"); -extern __device__ __attribute__((const)) unsigned int __mbcnt_hi(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.hi"); - __device__ static inline unsigned int __lane_id() { return __mbcnt_hi(-1, __mbcnt_lo(-1, 0)); } +/* +HIP specific device functions +*/ + +// utility union type +union __u { + int i; + unsigned int u; + float f; +}; + +__device__ static inline unsigned __hip_ds_bpermute(int index, unsigned src) { + __u tmp; tmp.u = src; + tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i); + return tmp.u; +} + +__device__ static inline float __hip_ds_bpermutef(int index, float src) { + __u tmp; tmp.f = src; + tmp.i = __llvm_amdgcn_ds_bpermute(index, tmp.i); + return tmp.f; +} + +__device__ static inline unsigned __hip_ds_permute(int index, unsigned src) { + __u tmp; tmp.u = src; + tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i); + return tmp.u; +} + +__device__ static inline float __hip_ds_permutef(int index, float src) { + __u tmp; tmp.u = src; + tmp.i = __llvm_amdgcn_ds_permute(index, tmp.i); + return tmp.u; +} + +__device__ static inline unsigned __hip_ds_swizzle(unsigned int src, int pattern) { + __u tmp; tmp.u = src; + tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern); + return tmp.u; +} +__device__ static inline float __hip_ds_swizzlef(float src, int pattern) { + __u tmp; tmp.f = src; + tmp.i = __llvm_amdgcn_ds_swizzle(tmp.i, pattern); + return tmp.f; +} + +__device__ static inline int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, + int bank_mask, bool bound_ctrl) { + return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); +} + +#define MASK1 0x00ff00ff +#define MASK2 0xff00ff00 + +__device__ static inline char4 __hip_hc_add8pk(char4 in1, char4 in2) { + char4 out; + unsigned one1 = in1.a & MASK1; + unsigned one2 = in2.a & MASK1; + out.a = (one1 + one2) & MASK1; + one1 = in1.a & MASK2; + one2 = in2.a & MASK2; + out.a = out.a | ((one1 + one2) & MASK2); + return out; +} + +__device__ static inline char4 __hip_hc_sub8pk(char4 in1, char4 in2) { + char4 out; + unsigned one1 = in1.a & MASK1; + unsigned one2 = in2.a & MASK1; + out.a = (one1 - one2) & MASK1; + one1 = in1.a & MASK2; + one2 = in2.a & MASK2; + out.a = out.a | ((one1 - one2) & MASK2); + return out; +} + +__device__ static inline char4 __hip_hc_mul8pk(char4 in1, char4 in2) { + char4 out; + unsigned one1 = in1.a & MASK1; + unsigned one2 = in2.a & MASK1; + out.a = (one1 * one2) & MASK1; + one1 = in1.a & MASK2; + one2 = in2.a & MASK2; + out.a = out.a | ((one1 * one2) & MASK2); + return out; +} + /* Rounding modes are not yet supported in HIP */ -__device__ float __double2float_rd(double x); -__device__ float __double2float_rn(double x); -__device__ float __double2float_ru(double x); -__device__ float __double2float_rz(double x); +__device__ static inline float __double2float_rd(double x) { return (double)x; } +__device__ static inline float __double2float_rn(double x) { return (double)x; } +__device__ static inline float __double2float_ru(double x) { return (double)x; } +__device__ static inline float __double2float_rz(double x) { return (double)x; } -__device__ int __double2hiint(double x); +__device__ static inline int __double2hiint(double x) { + static_assert(sizeof(double) == 2 * sizeof(int), ""); -__device__ int __double2int_rd(double x); -__device__ int __double2int_rn(double x); -__device__ int __double2int_ru(double x); -__device__ int __double2int_rz(double x); + int tmp[2]; + __builtin_memcpy(tmp, &x, sizeof(tmp)); -__device__ long long int __double2ll_rd(double x); -__device__ long long int __double2ll_rn(double x); -__device__ long long int __double2ll_ru(double x); -__device__ long long int __double2ll_rz(double x); + return tmp[1]; +} +__device__ static inline int __double2loint(double x) { + static_assert(sizeof(double) == 2 * sizeof(int), ""); -__device__ int __double2loint(double x); + int tmp[2]; + __builtin_memcpy(tmp, &x, sizeof(tmp)); -__device__ unsigned int __double2uint_rd(double x); -__device__ unsigned int __double2uint_rn(double x); -__device__ unsigned int __double2uint_ru(double x); -__device__ unsigned int __double2uint_rz(double x); + return tmp[0]; +} -__device__ unsigned long long int __double2ull_rd(double x); -__device__ unsigned long long int __double2ull_rn(double x); -__device__ unsigned long long int __double2ull_ru(double x); -__device__ unsigned long long int __double2ull_rz(double x); +__device__ static inline int __double2int_rd(double x) { return (int)x; } +__device__ static inline int __double2int_rn(double x) { return (int)x; } +__device__ static inline int __double2int_ru(double x) { return (int)x; } +__device__ static inline int __double2int_rz(double x) { return (int)x; } + +__device__ static inline long long int __double2ll_rd(double x) { return (long long int)x; } +__device__ static inline long long int __double2ll_rn(double x) { return (long long int)x; } +__device__ static inline long long int __double2ll_ru(double x) { return (long long int)x; } +__device__ static inline long long int __double2ll_rz(double x) { return (long long int)x; } + +__device__ static inline unsigned int __double2uint_rd(double x) { return (unsigned int)x; } +__device__ static inline unsigned int __double2uint_rn(double x) { return (unsigned int)x; } +__device__ static inline unsigned int __double2uint_ru(double x) { return (unsigned int)x; } +__device__ static inline unsigned int __double2uint_rz(double x) { return (unsigned int)x; } + +__device__ static inline unsigned long long int __double2ull_rd(double x) { + return (unsigned long long int)x; +} +__device__ static inline unsigned long long int __double2ull_rn(double x) { + return (unsigned long long int)x; +} +__device__ static inline unsigned long long int __double2ull_ru(double x) { + return (unsigned long long int)x; +} +__device__ static inline unsigned long long int __double2ull_rz(double x) { + return (unsigned long long int)x; +} + +__device__ static inline long long int __double_as_longlong(double x) { + static_assert(sizeof(long long) == sizeof(double), ""); + + long long tmp; + __builtin_memcpy(&tmp, &x, sizeof(tmp)); + + return tmp; +} -__device__ long long int __double_as_longlong(double x); /* __device__ unsigned short __float2half_rn(float x); __device__ float __half2float(unsigned short); @@ -146,72 +395,122 @@ CUDA implements half as unsigned short whereas, HIP doesn't. */ -__device__ int __float2int_rd(float x); -__device__ int __float2int_rn(float x); -__device__ int __float2int_ru(float x); -__device__ int __float2int_rz(float x); +__device__ static inline int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); } +__device__ static inline int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); } +__device__ static inline int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); } +__device__ static inline int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); } -__device__ long long int __float2ll_rd(float x); -__device__ long long int __float2ll_rn(float x); -__device__ long long int __float2ll_ru(float x); -__device__ long long int __float2ll_rz(float x); +__device__ static inline long long int __float2ll_rd(float x) { return (long long int)x; } +__device__ static inline long long int __float2ll_rn(float x) { return (long long int)x; } +__device__ static inline long long int __float2ll_ru(float x) { return (long long int)x; } +__device__ static inline long long int __float2ll_rz(float x) { return (long long int)x; } -__device__ unsigned int __float2uint_rd(float x); -__device__ unsigned int __float2uint_rn(float x); -__device__ unsigned int __float2uint_ru(float x); -__device__ unsigned int __float2uint_rz(float x); +__device__ static inline unsigned int __float2uint_rd(float x) { return (unsigned int)x; } +__device__ static inline unsigned int __float2uint_rn(float x) { return (unsigned int)x; } +__device__ static inline unsigned int __float2uint_ru(float x) { return (unsigned int)x; } +__device__ static inline unsigned int __float2uint_rz(float x) { return (unsigned int)x; } -__device__ unsigned long long int __float2ull_rd(float x); -__device__ unsigned long long int __float2ull_rn(float x); -__device__ unsigned long long int __float2ull_ru(float x); -__device__ unsigned long long int __float2ull_rz(float x); +__device__ static inline unsigned long long int __float2ull_rd(float x) { + return (unsigned long long int)x; +} +__device__ static inline unsigned long long int __float2ull_rn(float x) { + return (unsigned long long int)x; +} +__device__ static inline unsigned long long int __float2ull_ru(float x) { + return (unsigned long long int)x; +} +__device__ static inline unsigned long long int __float2ull_rz(float x) { + return (unsigned long long int)x; +} -__device__ int __float_as_int(float x); -__device__ unsigned int __float_as_uint(float x); -__device__ double __hiloint2double(int hi, int lo); -__device__ double __int2double_rn(int x); +__device__ static inline int __float_as_int(float x) { + static_assert(sizeof(int) == sizeof(float), ""); -__device__ float __int2float_rd(int x); -__device__ float __int2float_rn(int x); -__device__ float __int2float_ru(int x); -__device__ float __int2float_rz(int x); + int tmp; + __builtin_memcpy(&tmp, &x, sizeof(tmp)); -__device__ float __int_as_float(int x); + return tmp; +} -__device__ double __ll2double_rd(long long int x); -__device__ double __ll2double_rn(long long int x); -__device__ double __ll2double_ru(long long int x); -__device__ double __ll2double_rz(long long int x); +__device__ static inline unsigned int __float_as_uint(float x) { + static_assert(sizeof(unsigned int) == sizeof(float), ""); -__device__ float __ll2float_rd(long long int x); -__device__ float __ll2float_rn(long long int x); -__device__ float __ll2float_ru(long long int x); -__device__ float __ll2float_rz(long long int x); + unsigned int tmp; + __builtin_memcpy(&tmp, &x, sizeof(tmp)); -__device__ double __longlong_as_double(long long int x); + return tmp; +} -__device__ double __uint2double_rn(int x); +__device__ static inline double __hiloint2double(int hi, int lo) { + static_assert(sizeof(double) == sizeof(uint64_t), ""); -__device__ float __uint2float_rd(unsigned int x); -__device__ float __uint2float_rn(unsigned int x); -__device__ float __uint2float_ru(unsigned int x); -__device__ float __uint2float_rz(unsigned int x); + uint64_t tmp0 = (static_cast(hi) << 32ull) | static_cast(lo); + double tmp1; + __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); -__device__ float __uint_as_float(unsigned int x); + return tmp1; +} -__device__ double __ull2double_rd(unsigned long long int x); -__device__ double __ull2double_rn(unsigned long long int x); -__device__ double __ull2double_ru(unsigned long long int x); -__device__ double __ull2double_rz(unsigned long long int x); +__device__ static inline double __int2double_rn(int x) { return (double)x; } -__device__ float __ull2float_rd(unsigned long long int x); -__device__ float __ull2float_rn(unsigned long long int x); -__device__ float __ull2float_ru(unsigned long long int x); -__device__ float __ull2float_rz(unsigned long long int x); +__device__ static inline float __int2float_rd(int x) { return (float)x; } +__device__ static inline float __int2float_rn(int x) { return (float)x; } +__device__ static inline float __int2float_ru(int x) { return (float)x; } +__device__ static inline float __int2float_rz(int x) { return (float)x; } -__device__ char4 __hip_hc_add8pk(char4, char4); -__device__ char4 __hip_hc_sub8pk(char4, char4); -__device__ char4 __hip_hc_mul8pk(char4, char4); +__device__ static inline float __int_as_float(int x) { + static_assert(sizeof(float) == sizeof(int), ""); + + float tmp; + __builtin_memcpy(&tmp, &x, sizeof(tmp)); + + return tmp; +} + +__device__ static inline double __ll2double_rd(long long int x) { return (double)x; } +__device__ static inline double __ll2double_rn(long long int x) { return (double)x; } +__device__ static inline double __ll2double_ru(long long int x) { return (double)x; } +__device__ static inline double __ll2double_rz(long long int x) { return (double)x; } + +__device__ static inline float __ll2float_rd(long long int x) { return (float)x; } +__device__ static inline float __ll2float_rn(long long int x) { return (float)x; } +__device__ static inline float __ll2float_ru(long long int x) { return (float)x; } +__device__ static inline float __ll2float_rz(long long int x) { return (float)x; } + +__device__ static inline double __longlong_as_double(long long int x) { + static_assert(sizeof(double) == sizeof(long long), ""); + + double tmp; + __builtin_memcpy(&tmp, &x, sizeof(tmp)); + + return x; +} + +__device__ static inline double __uint2double_rn(int x) { return (double)x; } + +__device__ static inline float __uint2float_rd(unsigned int x) { return (float)x; } +__device__ static inline float __uint2float_rn(unsigned int x) { return (float)x; } +__device__ static inline float __uint2float_ru(unsigned int x) { return (float)x; } +__device__ static inline float __uint2float_rz(unsigned int x) { return (float)x; } + +__device__ static inline float __uint_as_float(unsigned int x) { + static_assert(sizeof(float) == sizeof(unsigned int), ""); + + float tmp; + __builtin_memcpy(&tmp, &x, sizeof(tmp)); + + return tmp; +} + +__device__ static inline double __ull2double_rd(unsigned long long int x) { return (double)x; } +__device__ static inline double __ull2double_rn(unsigned long long int x) { return (double)x; } +__device__ static inline double __ull2double_ru(unsigned long long int x) { return (double)x; } +__device__ static inline double __ull2double_rz(unsigned long long int x) { return (double)x; } + +__device__ static inline float __ull2float_rd(unsigned long long int x) { return (float)x; } +__device__ static inline float __ull2float_rn(unsigned long long int x) { return (float)x; } +__device__ static inline float __ull2float_ru(unsigned long long int x) { return (float)x; } +__device__ static inline float __ull2float_rz(unsigned long long int x) { return (float)x; } #if defined(__HCC__) #define __HCC_OR_HIP_CLANG__ 1 diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index b3fa556bd4..8bf3ce3a19 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -32,6 +32,11 @@ THE SOFTWARE. extern "C" __device__ int32_t __ockl_activelane_u32(void); +extern "C" __device__ float __ocml_floor_f32(float); +extern "C" __device__ float __ocml_rint_f32(float); +extern "C" __device__ float __ocml_ceil_f32(float); +extern "C" __device__ float __ocml_trunc_f32(float); + // Introduce local address space #define __local __attribute__((address_space(3))) __device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } diff --git a/include/hip/hcc_detail/llvm_intrinsics.h b/include/hip/hcc_detail/llvm_intrinsics.h index 2c7819b535..02df3c2fbe 100644 --- a/include/hip/hcc_detail/llvm_intrinsics.h +++ b/include/hip/hcc_detail/llvm_intrinsics.h @@ -34,4 +34,33 @@ THE SOFTWARE. __device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); +__device__ +unsigned int __llvm_bitrev_b32(unsigned int src0) __asm("llvm.bitreverse.i32"); + +__device__ +uint64_t __llvm_bitrev_b64(uint64_t src0) __asm("llvm.bitreverse.i64"); + +extern +__device__ +__attribute__((const)) +unsigned int __mbcnt_lo(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.lo"); + +extern +__device__ +__attribute__((const)) +unsigned int __mbcnt_hi(unsigned int x, unsigned int y) __asm("llvm.amdgcn.mbcnt.hi"); + +__device__ +int __llvm_amdgcn_ds_bpermute(int index, int src) __asm("llvm.amdgcn.ds.bpermute"); + +__device__ +int __llvm_amdgcn_ds_permute(int index, int src) __asm("llvm.amdgcn.ds.permute"); + +__device__ +int __llvm_amdgcn_ds_swizzle(int index, int pattern) __asm("llvm.amdgcn.ds.swizzle"); + +__device__ +int __llvm_amdgcn_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, + bool bound_ctrl) __asm("llvm.amdgcn.mov.dpp.i32"); + #endif diff --git a/src/device_functions.cpp b/src/device_functions.cpp deleted file mode 100644 index 8ef19bab3f..0000000000 --- a/src/device_functions.cpp +++ /dev/null @@ -1,387 +0,0 @@ -/* -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. -*/ - -#include -#include -#include -#include -#include "device_util.h" - -__device__ float __double2float_rd(double x) { return (double)x; } -__device__ float __double2float_rn(double x) { return (double)x; } -__device__ float __double2float_ru(double x) { return (double)x; } -__device__ float __double2float_rz(double x) { return (double)x; } - - -__device__ int __double2hiint(double x) { - static_assert(sizeof(double) == 2 * sizeof(int), ""); - - int tmp[2]; - __builtin_memcpy(tmp, &x, sizeof(tmp)); - - return tmp[1]; -} -__device__ int __double2loint(double x) { - static_assert(sizeof(double) == 2 * sizeof(int), ""); - - int tmp[2]; - __builtin_memcpy(tmp, &x, sizeof(tmp)); - - return tmp[0]; -} - - -__device__ int __double2int_rd(double x) { return (int)x; } -__device__ int __double2int_rn(double x) { return (int)x; } -__device__ int __double2int_ru(double x) { return (int)x; } -__device__ int __double2int_rz(double x) { return (int)x; } - -__device__ long long int __double2ll_rd(double x) { return (long long int)x; } -__device__ long long int __double2ll_rn(double x) { return (long long int)x; } -__device__ long long int __double2ll_ru(double x) { return (long long int)x; } -__device__ long long int __double2ll_rz(double x) { return (long long int)x; } - - -__device__ unsigned int __double2uint_rd(double x) { return (unsigned int)x; } -__device__ unsigned int __double2uint_rn(double x) { return (unsigned int)x; } -__device__ unsigned int __double2uint_ru(double x) { return (unsigned int)x; } -__device__ unsigned int __double2uint_rz(double x) { return (unsigned int)x; } - -__device__ unsigned long long int __double2ull_rd(double x) { return (unsigned long long int)x; } -__device__ unsigned long long int __double2ull_rn(double x) { return (unsigned long long int)x; } -__device__ unsigned long long int __double2ull_ru(double x) { return (unsigned long long int)x; } -__device__ unsigned long long int __double2ull_rz(double x) { return (unsigned long long int)x; } - -__device__ long long int __double_as_longlong(double x) { - static_assert(sizeof(long long) == sizeof(double), ""); - - long long tmp; - __builtin_memcpy(&tmp, &x, sizeof(tmp)); - - return tmp; -} - -__device__ int __float2int_rd(float x) { return (int)__ocml_floor_f32(x); } -__device__ int __float2int_rn(float x) { return (int)__ocml_rint_f32(x); } -__device__ int __float2int_ru(float x) { return (int)__ocml_ceil_f32(x); } -__device__ int __float2int_rz(float x) { return (int)__ocml_trunc_f32(x); } - -__device__ long long int __float2ll_rd(float x) { return (long long int)x; } -__device__ long long int __float2ll_rn(float x) { return (long long int)x; } -__device__ long long int __float2ll_ru(float x) { return (long long int)x; } -__device__ long long int __float2ll_rz(float x) { return (long long int)x; } - -__device__ unsigned int __float2uint_rd(float x) { return (unsigned int)x; } -__device__ unsigned int __float2uint_rn(float x) { return (unsigned int)x; } -__device__ unsigned int __float2uint_ru(float x) { return (unsigned int)x; } -__device__ unsigned int __float2uint_rz(float x) { return (unsigned int)x; } - -__device__ unsigned long long int __float2ull_rd(float x) { return (unsigned long long int)x; } -__device__ unsigned long long int __float2ull_rn(float x) { return (unsigned long long int)x; } -__device__ unsigned long long int __float2ull_ru(float x) { return (unsigned long long int)x; } -__device__ unsigned long long int __float2ull_rz(float x) { return (unsigned long long int)x; } - -__device__ int __float_as_int(float x) { - static_assert(sizeof(int) == sizeof(float), ""); - - int tmp; - __builtin_memcpy(&tmp, &x, sizeof(tmp)); - - return tmp; -} -__device__ unsigned int __float_as_uint(float x) { - static_assert(sizeof(unsigned int) == sizeof(float), ""); - - unsigned int tmp; - __builtin_memcpy(&tmp, &x, sizeof(tmp)); - - return tmp; -} -__device__ double __hiloint2double(int32_t hi, int32_t lo) { - static_assert(sizeof(double) == sizeof(uint64_t), ""); - - uint64_t tmp0 = (static_cast(hi) << 32ull) | static_cast(lo); - double tmp1; - __builtin_memcpy(&tmp1, &tmp0, sizeof(tmp0)); - - return tmp1; -} -__device__ double __int2double_rn(int x) { return (double)x; } - -__device__ float __int2float_rd(int x) { return (float)x; } -__device__ float __int2float_rn(int x) { return (float)x; } -__device__ float __int2float_ru(int x) { return (float)x; } -__device__ float __int2float_rz(int x) { return (float)x; } - -__device__ float __int_as_float(int x) { - static_assert(sizeof(float) == sizeof(int), ""); - - float tmp; - __builtin_memcpy(&tmp, &x, sizeof(tmp)); - - return tmp; -} - -__device__ double __ll2double_rd(long long int x) { return (double)x; } -__device__ double __ll2double_rn(long long int x) { return (double)x; } -__device__ double __ll2double_ru(long long int x) { return (double)x; } -__device__ double __ll2double_rz(long long int x) { return (double)x; } - -__device__ float __ll2float_rd(long long int x) { return (float)x; } -__device__ float __ll2float_rn(long long int x) { return (float)x; } -__device__ float __ll2float_ru(long long int x) { return (float)x; } -__device__ float __ll2float_rz(long long int x) { return (float)x; } - -__device__ double __longlong_as_double(long long int x) { - static_assert(sizeof(double) == sizeof(long long), ""); - - double tmp; - __builtin_memcpy(&tmp, &x, sizeof(tmp)); - - return x; -} - -__device__ double __uint2double_rn(int x) { return (double)x; } - -__device__ float __uint2float_rd(unsigned int x) { return (float)x; } -__device__ float __uint2float_rn(unsigned int x) { return (float)x; } -__device__ float __uint2float_ru(unsigned int x) { return (float)x; } -__device__ float __uint2float_rz(unsigned int x) { return (float)x; } - -__device__ float __uint_as_float(unsigned int x) { - static_assert(sizeof(float) == sizeof(unsigned int), ""); - - float tmp; - __builtin_memcpy(&tmp, &x, sizeof(tmp)); - - return tmp; -} - -__device__ double __ull2double_rd(unsigned long long int x) { return (double)x; } -__device__ double __ull2double_rn(unsigned long long int x) { return (double)x; } -__device__ double __ull2double_ru(unsigned long long int x) { return (double)x; } -__device__ double __ull2double_rz(unsigned long long int x) { return (double)x; } - -__device__ float __ull2float_rd(unsigned long long int x) { return (float)x; } -__device__ float __ull2float_rn(unsigned long long int x) { return (float)x; } -__device__ float __ull2float_ru(unsigned long long int x) { return (float)x; } -__device__ float __ull2float_rz(unsigned long long int x) { return (float)x; } - -/* -Integer Intrinsics -*/ - -// integer intrinsic function __poc __clz __ffs __brev -__device__ unsigned int __popc(unsigned int input) { return hc::__popcount_u32_b32(input); } - -__device__ unsigned int __popcll(unsigned long long int input) { - return hc::__popcount_u32_b64(input); -} - -__device__ unsigned int __clz(unsigned int input) { -#ifdef NVCC_COMPAT - return input == 0 ? 32 : hc::__firstbit_u32_u32(input); -#else - return hc::__firstbit_u32_u32(input); -#endif -} - -__device__ unsigned int __clzll(unsigned long long int input) { -#ifdef NVCC_COMPAT - return input == 0 ? 64 : hc::__firstbit_u32_u64(input); -#else - return hc::__firstbit_u32_u64(input); -#endif -} - -__device__ unsigned int __clz(int input) { -#ifdef NVCC_COMPAT - return input == 0 ? 32 : hc::__firstbit_u32_s32(input); -#else - return hc::__firstbit_u32_s32(input); -#endif -} - -__device__ unsigned int __clzll(long long int input) { -#ifdef NVCC_COMPAT - return input == 0 ? 64 : hc::__firstbit_u32_s64(input); -#else - return hc::__firstbit_u32_s64(input); -#endif -} - -__device__ unsigned int __ffs(unsigned int input) { -#ifdef NVCC_COMPAT - return hc::__lastbit_u32_u32(input) + 1; -#else - return hc::__lastbit_u32_u32(input); -#endif -} - -__device__ unsigned int __ffsll(unsigned long long int input) { -#ifdef NVCC_COMPAT - return hc::__lastbit_u32_u64(input) + 1; -#else - return hc::__lastbit_u32_u64(input); -#endif -} - -__device__ unsigned int __ffs(int input) { -#ifdef NVCC_COMPAT - return hc::__lastbit_u32_s32(input) + 1; -#else - return hc::__lastbit_u32_s32(input); -#endif -} - -__device__ unsigned int __ffsll(long long int input) { -#ifdef NVCC_COMPAT - return hc::__lastbit_u32_s64(input) + 1; -#else - return hc::__lastbit_u32_s64(input); -#endif -} - -__device__ unsigned int __brev(unsigned int input) { return hc::__bitrev_b32(input); } - -__device__ unsigned long long int __brevll(unsigned long long int input) { - return hc::__bitrev_b64(input); -} - -struct ucharHolder { - union { - unsigned char c[4]; - unsigned int ui; - }; -} __attribute__((aligned(4))); - -struct uchar2Holder { - union { - unsigned int ui[2]; - unsigned char c[8]; - }; -} __attribute__((aligned(8))); - -__device__ unsigned int __byte_perm(unsigned int x, unsigned int y, unsigned int s) { - struct uchar2Holder cHoldVal; - struct ucharHolder cHoldKey; - struct ucharHolder cHoldOut; - cHoldKey.ui = s; - cHoldVal.ui[0] = x; - cHoldVal.ui[1] = y; - cHoldOut.c[0] = cHoldVal.c[cHoldKey.c[0]]; - cHoldOut.c[1] = cHoldVal.c[cHoldKey.c[1]]; - cHoldOut.c[2] = cHoldVal.c[cHoldKey.c[2]]; - cHoldOut.c[3] = cHoldVal.c[cHoldKey.c[3]]; - return cHoldOut.ui; -} - -__device__ long long __mul64hi(long long int x, long long int y) { - ulong x0 = (ulong)x & 0xffffffffUL; - long x1 = x >> 32; - ulong y0 = (ulong)y & 0xffffffffUL; - long y1 = y >> 32; - ulong z0 = x0*y0; - long t = x1*y0 + (z0 >> 32); - long z1 = t & 0xffffffffL; - long z2 = t >> 32; - z1 = x0*y1 + z1; - return x1*y1 + z2 + (z1 >> 32); -} - -__device__ unsigned long long __umul64hi(unsigned long long int x, unsigned long long int y) { - ulong x0 = x & 0xffffffffUL; - ulong x1 = x >> 32; - ulong y0 = y & 0xffffffffUL; - ulong y1 = y >> 32; - ulong z0 = x0*y0; - ulong t = x1*y0 + (z0 >> 32); - ulong z1 = t & 0xffffffffUL; - ulong z2 = t >> 32; - z1 = x0*y1 + z1; - return x1*y1 + z2 + (z1 >> 32); -} - -/* -HIP specific device functions -*/ - -__device__ unsigned __hip_ds_bpermute(int index, unsigned src) { - return hc::__amdgcn_ds_bpermute(index, src); -} - -__device__ float __hip_ds_bpermutef(int index, float src) { - return hc::__amdgcn_ds_bpermute(index, src); -} - -__device__ unsigned __hip_ds_permute(int index, unsigned src) { - return hc::__amdgcn_ds_permute(index, src); -} - -__device__ float __hip_ds_permutef(int index, float src) { - return hc::__amdgcn_ds_permute(index, src); -} - -__device__ unsigned __hip_ds_swizzle(unsigned int src, int pattern) { - return hc::__amdgcn_ds_swizzle(src, pattern); -} - -__device__ float __hip_ds_swizzlef(float src, int pattern) { - return hc::__amdgcn_ds_swizzle(src, pattern); -} - -__device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask, bool bound_ctrl) { - return hc::__amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); -} - -#define MASK1 0x00ff00ff -#define MASK2 0xff00ff00 - -__device__ char4 __hip_hc_add8pk(char4 in1, char4 in2) { - char4 out; - unsigned one1 = in1.a & MASK1; - unsigned one2 = in2.a & MASK1; - out.a = (one1 + one2) & MASK1; - one1 = in1.a & MASK2; - one2 = in2.a & MASK2; - out.a = out.a | ((one1 + one2) & MASK2); - return out; -} - -__device__ char4 __hip_hc_sub8pk(char4 in1, char4 in2) { - char4 out; - unsigned one1 = in1.a & MASK1; - unsigned one2 = in2.a & MASK1; - out.a = (one1 - one2) & MASK1; - one1 = in1.a & MASK2; - one2 = in2.a & MASK2; - out.a = out.a | ((one1 - one2) & MASK2); - return out; -} - -__device__ char4 __hip_hc_mul8pk(char4 in1, char4 in2) { - char4 out; - unsigned one1 = in1.a & MASK1; - unsigned one2 = in2.a & MASK1; - out.a = (one1 * one2) & MASK1; - one1 = in1.a & MASK2; - one2 = in2.a & MASK2; - out.a = out.a | ((one1 * one2) & MASK2); - return out; -} From 2142eb4d1263e1dbe6183f4656db3416d0fb7fb0 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 19 Jun 2018 21:09:44 +0000 Subject: [PATCH 08/21] Implement hip_hc.ll into HIP headers Move all __hip_hc_ir_* functions from hip_hc.ll into HIP header as inline asm. Remove hip_hc.ll and build dependencies from HIP. --- CMakeLists.txt | 1 - bin/hipcc | 5 +- include/hip/hcc_detail/device_functions.h | 56 ++++++++++++++++++++--- packaging/hip_hcc.txt | 1 - src/hip_hc.ll | 30 ------------ 5 files changed, 52 insertions(+), 41 deletions(-) delete mode 100644 src/hip_hc.ll diff --git a/CMakeLists.txt b/CMakeLists.txt index b649b20c21..9095ff6531 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -237,7 +237,6 @@ endif() # Install hip_hcc if platform is hcc if(HIP_PLATFORM STREQUAL "hcc") install(TARGETS hip_hcc_static hip_hcc hip_device DESTINATION lib) - install(FILES ${CMAKE_CURRENT_SOURCE_DIR}/src/hip_hc.ll DESTINATION lib) # Install .hipInfo install(FILES ${PROJECT_BINARY_DIR}/.hipInfo DESTINATION lib) diff --git a/bin/hipcc b/bin/hipcc index 03f35b27fc..4f56ffd875 100755 --- a/bin/hipcc +++ b/bin/hipcc @@ -254,8 +254,7 @@ if($HIP_PLATFORM eq "hcc"){ } if(($HIP_PLATFORM eq "hcc")){ - $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n"; - $ENV{HIP_HC_IR_FILE}=""; + $ENV{HCC_EXTRA_LIBRARIES}="\n"; } if($HIP_PLATFORM eq "nvcc"){ @@ -508,7 +507,7 @@ if($HIP_PLATFORM eq "hcc" or $HIP_PLATFORM eq "clang"){ print "No valid AMD GPU target was either specified or found. Please specify a valid target using --amdgpu-target=" and die(); } - $ENV{HCC_EXTRA_LIBRARIES}="$HIP_PATH/lib/hip_hc.ll\n"; + $ENV{HCC_EXTRA_LIBRARIES}="\n"; if($HIP_PLATFORM eq "hcc") { $GPU_ARCH_OPT = " --amdgpu-target="; diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index b147cd9b80..cb9dd82c0d 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -34,11 +34,55 @@ THE SOFTWARE. typedef unsigned long ulong; typedef unsigned int uint; -extern "C" __device__ unsigned int __hip_hc_ir_umul24_int(unsigned int, unsigned int); -extern "C" __device__ signed int __hip_hc_ir_mul24_int(signed int, signed int); -extern "C" __device__ signed int __hip_hc_ir_mulhi_int(signed int, signed int); -extern "C" __device__ unsigned int __hip_hc_ir_umulhi_int(unsigned int, unsigned int); -extern "C" __device__ unsigned int __hip_hc_ir_usad_int(unsigned int, unsigned int, unsigned int); +extern "C" __device__ inline uint __hip_hc_ir_umul24_int(uint a, uint b) { + // define i32 @__hip_hc_ir_umul24_int(i32 %a, i32 %b) #1 { + // %1 = tail call i32 asm sideeffect "v_mul_u32_u24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) + // ret i32 %1 + // } + uint out; + __asm volatile("v_mul_u32_u24 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); + return out; +} + +extern "C" __device__ inline int __hip_hc_ir_mul24_int(int a, int b) { + // define i32 @__hip_hc_ir_mul24_int(i32 %a, i32 %b) #1 { + // %1 = tail call i32 asm sideeffect "v_mul_i32_i24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) + // ret i32 %1 + // } + int out; + __asm volatile("v_mul_i32_i24 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); + return out; +} + +extern "C" __device__ inline int __hip_hc_ir_mulhi_int(int a, int b) { + // define i32 @__hip_hc_ir_mulhi_int(i32 %a, i32 %b) #1 { + // %1 = tail call i32 asm sideeffect "v_mul_hi_i32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) + // ret i32 %1 + // } + int out; + __asm volatile("v_mul_hi_i32 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); + return out; +} + +extern "C" __device__ inline uint __hip_hc_ir_umulhi_int(uint a, uint b) { + // define i32 @__hip_hc_ir_umulhi_int(i32 %a, i32 %b) #1 { + // %1 = tail call i32 asm sideeffect "v_mul_hi_u32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) + // ret i32 %1 + // } + uint out; + __asm volatile("v_mul_hi_u32 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); + return out; +} + +extern "C" __device__ inline uint __hip_hc_ir_usad_int(uint a, uint b, uint c) { + // define i32 @__hip_hc_ir_usad_int(i32 %a, i32 %b, i32 %c) #1 { + // %1 = tail call i32 asm sideeffect "v_sad_u32 $0, $1, $2, $3","=v,v,v,v"(i32 %a, i32 %b, i32 %c) + // ret i32 %1 + // } + uint out; + __asm volatile("v_sad_u32 %0, %1, %2, %3" : "=v"(out) : "v"(a), "v"(b), "v"(c)); + return out; +} /* Integer Intrinsics @@ -556,7 +600,7 @@ uint64_t __ballot64(int a) { // %b = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %a) #1 // ret i64 %b // } - __asm("v_cmp_ne_i32_e64 $0, 0, $1" : "=s"(s) : "v"(a)); + __asm("v_cmp_ne_i32_e64 %0, 0, %1" : "=s"(s) : "v"(a)); return s; } diff --git a/packaging/hip_hcc.txt b/packaging/hip_hcc.txt index 04293f2044..9d4b96761d 100644 --- a/packaging/hip_hcc.txt +++ b/packaging/hip_hcc.txt @@ -5,7 +5,6 @@ install(FILES @PROJECT_BINARY_DIR@/libhip_hcc.so DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/libhip_hcc_static.a DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/libhip_device.a DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/.hipInfo DESTINATION lib) -install(FILES @hip_SOURCE_DIR@/src/hip_hc.ll DESTINATION lib) install(FILES @PROJECT_BINARY_DIR@/hip-config.cmake @PROJECT_BINARY_DIR@/hip-config-version.cmake DESTINATION lib/cmake/hip) install(FILES @hip_SOURCE_DIR@/packaging/hip-targets.cmake @hip_SOURCE_DIR@/packaging/hip-targets-release.cmake DESTINATION lib/cmake/hip) diff --git a/src/hip_hc.ll b/src/hip_hc.ll deleted file mode 100644 index aba9205912..0000000000 --- a/src/hip_hc.ll +++ /dev/null @@ -1,30 +0,0 @@ -target datalayout = "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64" -target triple = "amdgcn--amdhsa" - -define i32 @__hip_hc_ir_mul24_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_mul_i32_i24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - ret i32 %1 -} - -define i32 @__hip_hc_ir_umul24_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_mul_u32_u24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - ret i32 %1 -} - -define i32 @__hip_hc_ir_mulhi_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_mul_hi_i32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - ret i32 %1 -} - -define i32 @__hip_hc_ir_umulhi_int(i32 %a, i32 %b) #1 { - %1 = tail call i32 asm sideeffect "v_mul_hi_u32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - ret i32 %1 -} - -define i32 @__hip_hc_ir_usad_int(i32 %a, i32 %b, i32 %c) #1 { - %1 = tail call i32 asm sideeffect "v_sad_u32 $0, $1, $2, $3","=v,v,v,v"(i32 %a, i32 %b, i32 %c) - ret i32 %1 -} - -attributes #1 = { alwaysinline nounwind } - From 6dc16bbf04ab1f4310b714c2bf87097e761884c4 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 19 Jun 2018 22:07:26 +0000 Subject: [PATCH 09/21] Implement __ballot, __any, __all into HIP headers --- include/hip/hcc_detail/device_functions.h | 56 +++++++++++++++++------ src/device_util.cpp | 17 ------- 2 files changed, 42 insertions(+), 31 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index cb9dd82c0d..71963a99e0 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -585,23 +585,51 @@ void __named_sync(int a, int b) { __builtin_amdgcn_s_barrier(); } #endif // __HIP_DEVICE_COMPILE__ // warp vote function __all __any __ballot -__device__ -int __all(int input); -__device__ -int __any(int input); -__device__ -unsigned long long int __ballot(int input); +extern "C" __device__ inline uint64_t __activelanemask_v4_b64_b1(unsigned int input) { + uint64_t output; + // define i64 @__activelanemask_v4_b64_b1(i32 %input) #5 { + // %a = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %input) #9 + // ret i64 %a + // } + __asm("v_cmp_ne_i32_e64 %0, 0, %1" : "=s"(output) : "v"(input)); + return output; +} __device__ inline -uint64_t __ballot64(int a) { - int64_t s; - // define i64 @__ballot64(i32 %a) #0 { - // %b = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %a) #1 - // ret i64 %b - // } - __asm("v_cmp_ne_i32_e64 %0, 0, %1" : "=s"(s) : "v"(a)); - return s; +unsigned int __activelanecount_u32_b1(unsigned int input) { + return __popcll(__activelanemask_v4_b64_b1(input)); +} + +__device__ +inline +int __all(int predicate) { + return __popcll(__activelanemask_v4_b64_b1(predicate)) == __activelanecount_u32_b1(1); +} + +__device__ +inline +int __any(int predicate) { +#ifdef NVCC_COMPAT + if (__popcll(__activelanemask_v4_b64_b1(predicate)) != 0) + return 1; + else + return 0; +#else + return __popcll(__activelanemask_v4_b64_b1(predicate)); +#endif +} + +__device__ +inline +unsigned long long int __ballot(int predicate) { + return __activelanemask_v4_b64_b1(predicate); +} + +__device__ +inline +unsigned long long int __ballot64(int predicate) { + return __activelanemask_v4_b64_b1(predicate); } // hip.amdgcn.bc - lanemask diff --git a/src/device_util.cpp b/src/device_util.cpp index a3386ba14d..853ca71c09 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -147,23 +147,6 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) { // abort __device__ void abort() { return hc::abort(); } -// warp vote function __all __any __ballot -__device__ int __all(int input) { return hc::__all(input); } - - -__device__ int __any(int input) { -#ifdef NVCC_COMPAT - if (hc::__any(input) != 0) - return 1; - else - return 0; -#else - return hc::__any(input); -#endif -} - -__device__ unsigned long long int __ballot(int input) { return hc::__ballot(input); } - // warp shuffle functions __device__ int __shfl(int input, int lane, int width) { return hc::__shfl(input, lane, width); } From 8ac864c2e377ca71284aca0071a526b439700f3f Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 20 Jun 2018 20:39:30 +0000 Subject: [PATCH 10/21] Replace __hip_hc_ir_ inline asm with __ockl_* functions --- include/hip/hcc_detail/device_functions.h | 66 ++++--------------- include/hip/hcc_detail/device_library_decls.h | 6 ++ 2 files changed, 17 insertions(+), 55 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 71963a99e0..8ea6632ffc 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -34,56 +34,6 @@ THE SOFTWARE. typedef unsigned long ulong; typedef unsigned int uint; -extern "C" __device__ inline uint __hip_hc_ir_umul24_int(uint a, uint b) { - // define i32 @__hip_hc_ir_umul24_int(i32 %a, i32 %b) #1 { - // %1 = tail call i32 asm sideeffect "v_mul_u32_u24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - // ret i32 %1 - // } - uint out; - __asm volatile("v_mul_u32_u24 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); - return out; -} - -extern "C" __device__ inline int __hip_hc_ir_mul24_int(int a, int b) { - // define i32 @__hip_hc_ir_mul24_int(i32 %a, i32 %b) #1 { - // %1 = tail call i32 asm sideeffect "v_mul_i32_i24 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - // ret i32 %1 - // } - int out; - __asm volatile("v_mul_i32_i24 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); - return out; -} - -extern "C" __device__ inline int __hip_hc_ir_mulhi_int(int a, int b) { - // define i32 @__hip_hc_ir_mulhi_int(i32 %a, i32 %b) #1 { - // %1 = tail call i32 asm sideeffect "v_mul_hi_i32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - // ret i32 %1 - // } - int out; - __asm volatile("v_mul_hi_i32 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); - return out; -} - -extern "C" __device__ inline uint __hip_hc_ir_umulhi_int(uint a, uint b) { - // define i32 @__hip_hc_ir_umulhi_int(i32 %a, i32 %b) #1 { - // %1 = tail call i32 asm sideeffect "v_mul_hi_u32 $0, $1, $2","=v,v,v"(i32 %a, i32 %b) - // ret i32 %1 - // } - uint out; - __asm volatile("v_mul_hi_u32 %0, %1, %2" : "=v"(out) : "v"(a), "v"(b)); - return out; -} - -extern "C" __device__ inline uint __hip_hc_ir_usad_int(uint a, uint b, uint c) { - // define i32 @__hip_hc_ir_usad_int(i32 %a, i32 %b, i32 %c) #1 { - // %1 = tail call i32 asm sideeffect "v_sad_u32 $0, $1, $2, $3","=v,v,v,v"(i32 %a, i32 %b, i32 %c) - // ret i32 %1 - // } - uint out; - __asm volatile("v_sad_u32 %0, %1, %2, %3" : "=v"(out) : "v"(a), "v"(b), "v"(c)); - return out; -} - /* Integer Intrinsics */ @@ -217,7 +167,10 @@ __device__ static inline unsigned int __hadd(int x, int y) { int value = z & 0x7FFFFFFF; return ((value) >> 1 || sign); } -__device__ static inline int __mul24(int x, int y) { return __hip_hc_ir_mul24_int(x, y); } + +__device__ static inline int __mul24(int x, int y) { + return __ockl_mul24_i32(x, y); +} __device__ static inline long long __mul64hi(long long int x, long long int y) { ulong x0 = (ulong)x & 0xffffffffUL; @@ -232,7 +185,10 @@ __device__ static inline long long __mul64hi(long long int x, long long int y) { return x1*y1 + z2 + (z1 >> 32); } -__device__ static inline int __mulhi(int x, int y) { return __hip_hc_ir_mulhi_int(x, y); } +__device__ static inline int __mulhi(int x, int y) { + return __ockl_mul_hi_i32(x, y); +} + __device__ static inline int __rhadd(int x, int y) { int z = x + y + 1; int sign = z & 0x8000000; @@ -246,7 +202,7 @@ __device__ static inline unsigned int __uhadd(unsigned int x, unsigned int y) { return (x + y) >> 1; } __device__ static inline int __umul24(unsigned int x, unsigned int y) { - return __hip_hc_ir_umul24_int(x, y); + return __ockl_mul24_u32(x, y); } __device__ @@ -264,13 +220,13 @@ static inline unsigned long long __umul64hi(unsigned long long int x, unsigned l } __device__ static inline unsigned int __umulhi(unsigned int x, unsigned int y) { - return __hip_hc_ir_umulhi_int(x, y); + return __ockl_mul_hi_u32(x, y); } __device__ static inline unsigned int __urhadd(unsigned int x, unsigned int y) { return (x + y + 1) >> 1; } __device__ static inline unsigned int __usad(unsigned int x, unsigned int y, unsigned int z) { - return __hip_hc_ir_usad_int(x, y, z); + return __ockl_sad_u32(x, y, z); } __device__ static inline unsigned int __lane_id() { return __mbcnt_hi(-1, __mbcnt_lo(-1, 0)); } diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 8bf3ce3a19..fba2d53e2e 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -32,6 +32,12 @@ THE SOFTWARE. extern "C" __device__ int32_t __ockl_activelane_u32(void); +extern "C" __device__ uint __ockl_mul24_u32(uint, uint); +extern "C" __device__ int __ockl_mul24_i32(int, int); +extern "C" __device__ uint __ockl_mul_hi_u32(uint, uint); +extern "C" __device__ int __ockl_mul_hi_i32(int, int); +extern "C" __device__ uint __ockl_sad_u32(uint, uint, uint); + extern "C" __device__ float __ocml_floor_f32(float); extern "C" __device__ float __ocml_rint_f32(float); extern "C" __device__ float __ocml_ceil_f32(float); From 9ac31e0bb65029b164e40f76af4cfabf75fa1141 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Fri, 22 Jun 2018 19:11:35 +0000 Subject: [PATCH 11/21] Implement __shfl_* funcs into HIP headers --- include/hip/hcc_detail/device_functions.h | 121 ++++++++++++++++++++++ include/hip/hcc_detail/hip_runtime.h | 25 ----- src/device_util.cpp | 29 ------ 3 files changed, 121 insertions(+), 54 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 8ea6632ffc..aae0706033 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -282,6 +282,127 @@ __device__ static inline int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, return __llvm_amdgcn_move_dpp(src, dpp_ctrl, row_mask, bank_mask, bound_ctrl); } +static constexpr int warpSize = 64; + + __device__ +inline +int __shfl(int var, int src_lane, int width = warpSize) { + int self = __lane_id(); + int index = src_lane + (self & ~(width-1)); + return __llvm_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl(unsigned int var, int src_lane, int width = warpSize) { + __u tmp; tmp.u = var; + tmp.i = __shfl(tmp.i, src_lane, width); + return tmp.u; +} +__device__ +inline +float __shfl(float var, int src_lane, int width = warpSize) { + __u tmp; tmp.f = var; + tmp.i = __shfl(tmp.i, src_lane, width); + return tmp.f; +} +__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; +} + + __device__ +inline +int __shfl_up(int var, unsigned int lane_delta, int width = warpSize) { + int self = __lane_id(); + int index = self - lane_delta; + index = (index < (self & ~(width-1)))?self:index; + return __llvm_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl_up(unsigned int var, unsigned int lane_delta, int width = warpSize) { + __u tmp; tmp.u = var; + tmp.i = __shfl_up(tmp.i, lane_delta, width); + return tmp.u; +} +__device__ +inline +float __shfl_up(float var, unsigned int lane_delta, int width = warpSize) { + __u tmp; tmp.f = var; + tmp.i = __shfl_up(tmp.i, lane_delta, width); + return tmp.f; +} +__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; +} + +__device__ +inline +int __shfl_down(int var, unsigned int lane_delta, int width = warpSize) { + int self = __lane_id(); + int index = self + lane_delta; + index = (int)((self&(width-1))+lane_delta) >= width?self:index; + return __llvm_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl_down(unsigned int var, unsigned int lane_delta, int width = warpSize) { + __u tmp; tmp.u = var; + tmp.i = __shfl_down(tmp.i, lane_delta, width); + return tmp.u; +} +__device__ +inline +float __shfl_down(float var, unsigned int lane_delta, int width = warpSize) { + __u tmp; tmp.f = var; + tmp.i = __shfl_down(tmp.i, lane_delta, width); + return tmp.f; +} +__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; +} + +__device__ +inline +int __shfl_xor(int var, int lane_mask, int width = warpSize) { + int self = __lane_id(); + int index = self^lane_mask; + index = index >= ((self+width)&~(width-1))?self:index; + return __llvm_amdgcn_ds_bpermute(index<<2, var); +} +__device__ +inline +unsigned int __shfl_xor(unsigned int var, int lane_mask, int width = warpSize) { + __u tmp; tmp.u = var; + tmp.i = __shfl_xor(tmp.i, lane_mask, width); + return tmp.u; +} +__device__ +inline +float __shfl_xor(float var, int lane_mask, int width = warpSize) { + __u tmp; tmp.f = var; + tmp.i = __shfl_xor(tmp.i, lane_mask, width); + return tmp.f; +} +__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; +} + #define MASK1 0x00ff00ff #define MASK2 0xff00ff00 diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 18b04daf77..8107f00a4e 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -184,36 +184,11 @@ extern int HIP_TRACE_API; #if __HCC_OR_HIP_CLANG__ -// TODO - hipify-clang - change to use the function call. -//#define warpSize hc::__wavesize() -static constexpr int warpSize = 64; - // abort __device__ void abort(); #if __HIP_ARCH_GFX701__ == 0 -// warp shuffle functions -#ifdef __cplusplus -__device__ int __shfl(int input, int lane, int width = warpSize); -__device__ int __shfl_up(int input, unsigned int lane_delta, int width = warpSize); -__device__ int __shfl_down(int input, unsigned int lane_delta, int width = warpSize); -__device__ int __shfl_xor(int input, int lane_mask, int width = warpSize); -__device__ float __shfl(float input, int lane, int width = warpSize); -__device__ float __shfl_up(float input, unsigned int lane_delta, int width = warpSize); -__device__ float __shfl_down(float input, unsigned int lane_delta, int width = warpSize); -__device__ float __shfl_xor(float input, int lane_mask, int width = warpSize); -#else -__device__ int __shfl(int input, int lane, int width); -__device__ int __shfl_up(int input, unsigned int lane_delta, int width); -__device__ int __shfl_down(int input, unsigned int lane_delta, int width); -__device__ int __shfl_xor(int input, int lane_mask, int width); -__device__ float __shfl(float input, int lane, int width); -__device__ float __shfl_up(float input, unsigned int lane_delta, int width); -__device__ float __shfl_down(float input, unsigned int lane_delta, int width); -__device__ float __shfl_xor(float input, int lane_mask, int width); -#endif //__cplusplus - __device__ unsigned __hip_ds_bpermute(int index, unsigned src); __device__ float __hip_ds_bpermutef(int index, float src); __device__ unsigned __hip_ds_permute(int index, unsigned src); diff --git a/src/device_util.cpp b/src/device_util.cpp index 853ca71c09..65ee5f4368 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -147,35 +147,6 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) { // abort __device__ void abort() { return hc::abort(); } -// warp shuffle functions -__device__ int __shfl(int input, int lane, int width) { return hc::__shfl(input, lane, width); } - -__device__ int __shfl_up(int input, unsigned int lane_delta, int width) { - return hc::__shfl_up(input, lane_delta, width); -} - -__device__ int __shfl_down(int input, unsigned int lane_delta, int width) { - return hc::__shfl_down(input, lane_delta, width); -} - -__device__ int __shfl_xor(int input, int lane_mask, int width) { - return hc::__shfl_xor(input, lane_mask, width); -} - -__device__ float __shfl(float input, int lane, int width) { return hc::__shfl(input, lane, width); } - -__device__ float __shfl_up(float input, unsigned int lane_delta, int width) { - return hc::__shfl_up(input, lane_delta, width); -} - -__device__ float __shfl_down(float input, unsigned int lane_delta, int width) { - return hc::__shfl_down(input, lane_delta, width); -} - -__device__ float __shfl_xor(float input, int lane_mask, int width) { - return hc::__shfl_xor(input, lane_mask, width); -} - __host__ __device__ int min(int arg1, int arg2) { return (int)(hc::precise_math::fmin((float)arg1, (float)arg2)); } From 07de5cb3342bfbc8142020179e6f3680938eb270 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 4 Jul 2018 17:56:50 +0000 Subject: [PATCH 12/21] Workaround cast warning of smaller integer type for __to_local For now, guard the __to_local function for device compile only since a local pointer should be same size as unsigned int on GPU compile. Also change to void* instead of char*. --- include/hip/hcc_detail/device_functions.h | 4 ++++ include/hip/hcc_detail/device_library_decls.h | 6 +++++- 2 files changed, 9 insertions(+), 1 deletion(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index aae0706033..32509ffffd 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -734,6 +734,8 @@ int64_t __lanemask_lt() return ballot; } +#ifdef __HIP_DEVICE_COMPILE__ + __device__ inline void* __get_dynamicgroupbaseptr() @@ -748,6 +750,8 @@ void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } +#endif // __HIP_DEVICE_COMPILE__ + #endif // __HCC_OR_HIP_CLANG__ #ifdef __HCC__ diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index fba2d53e2e..a636c2c950 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -45,7 +45,11 @@ extern "C" __device__ float __ocml_trunc_f32(float); // Introduce local address space #define __local __attribute__((address_space(3))) -__device__ inline static __local char* __to_local(unsigned x) { return (__local char*)x; } + +#ifdef __HIP_DEVICE_COMPILE__ +__device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; } +#endif //__HIP_DEVICE_COMPILE__ + extern "C" __device__ void* __local_to_generic(__local void* p); // __llvm_fence* functions from device-libs/irif/src/fence.ll From 930a16bccdd49337bd063c1c7ed30e021308c3a1 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Wed, 4 Jul 2018 23:13:35 +0000 Subject: [PATCH 13/21] Implement Memory Fence Functions in header Enabled __llvm_fence_* functions for seq_cst. --- include/hip/hcc_detail/device_functions.h | 136 +++++++++++++----- include/hip/hcc_detail/device_library_decls.h | 11 ++ include/hip/hcc_detail/hip_runtime.h | 75 ---------- src/device_util.cpp | 3 - src/device_util.h | 1 - 5 files changed, 109 insertions(+), 117 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 32509ffffd..716c51a887 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -752,6 +752,101 @@ void *__amdgcn_get_dynamicgroupbaseptr() { #endif // __HIP_DEVICE_COMPILE__ + +// 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; + } + } +} + +// Memory Fence Functions +__device__ +inline +static void __threadfence() +{ + __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_device); +} + +__device__ +inline +static void __threadfence_block() +{ + __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_work_group); +} + +__device__ +inline +static void __threadfence_system() +{ + __atomic_work_item_fence(0, __memory_order_seq_cst, __memory_scope_all_svm_devices); +} + #endif // __HCC_OR_HIP_CLANG__ #ifdef __HCC__ @@ -796,50 +891,14 @@ void __assertfail(const char * __assertion, __builtin_trap(); } -// 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 __work_group_barrier(__cl_mem_fence_flags flags, __memory_scope scope) { if (flags) { - switch(scope) { - case __memory_scope_work_item: break; - case __memory_scope_sub_group: __llvm_fence_rel_sg(); break; - case __memory_scope_work_group: __llvm_fence_rel_wg(); break; - case __memory_scope_device: __llvm_fence_rel_dev(); break; - case __memory_scope_all_svm_devices: __llvm_fence_rel_sys(); break; - } - //atomic_work_item_fence(flags, memory_order_release, scope); + __atomic_work_item_fence(flags, __memory_order_release, scope); __builtin_amdgcn_s_barrier(); - //atomic_work_item_fence(flags, memory_order_acquire, scope); - switch(scope) { - case __memory_scope_work_item: break; - case __memory_scope_sub_group: __llvm_fence_acq_sg(); break; - case __memory_scope_work_group: __llvm_fence_acq_wg(); break; - case __memory_scope_device: __llvm_fence_acq_dev(); break; - case __memory_scope_all_svm_devices: __llvm_fence_acq_sys(); break; - } + __atomic_work_item_fence(flags, __memory_order_acquire, scope); } else { __builtin_amdgcn_s_barrier(); } @@ -918,4 +977,5 @@ unsigned __smid(void) #endif //defined(__clang__) && defined(__HIP__) + #endif diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index a636c2c950..82c39b24f0 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -63,4 +63,15 @@ 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); + #endif diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 68f2244014..3e0f2e27c5 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -204,81 +204,6 @@ __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); - -/** - * CUDA 8 device function features - - */ - - -/** - * Kernel launching - */ - -/** - *------------------------------------------------------------------------------------------------- - *------------------------------------------------------------------------------------------------- - * @defgroup Fence Fence Functions - * @{ - * - * - * @warning The HIP memory fence functions are currently not supported yet. - * If any of those threadfence stubs are reached by the application, you should set "export - *HSA_DISABLE_CACHE=1" to disable L1 and L2 caches. - * - * - * On AMD platforms, the threadfence* routines are currently empty stubs. - */ - -extern __attribute__((const)) __device__ void __hip_hc_threadfence() __asm("__llvm_fence_sc_dev"); -extern __attribute__((const)) __device__ void __hip_hc_threadfence_block() __asm( - "__llvm_fence_sc_wg"); - - -/** - * @brief threadfence_block makes writes visible to threads running in same block. - * - * @Returns void - * - * @param void - * - * @warning __threadfence_block is a stub and map to no-op. - */ -// __device__ void __threadfence_block(void); -__device__ static inline void __threadfence_block(void) { return __hip_hc_threadfence_block(); } - -/** - * @brief threadfence makes wirtes visible to other threads running on same GPU. - * - * @Returns void - * - * @param void - * - * @warning __threadfence is a stub and map to no-op, application should set "export - * HSA_DISABLE_CACHE=1" to disable both L1 and L2 caches. - */ -// __device__ void __threadfence(void) __attribute__((deprecated("Provided for compile-time -// compatibility, not yet functional"))); -__device__ static inline void __threadfence(void) { return __hip_hc_threadfence(); } - -/** - * @brief threadfence_system makes writes to pinned system memory visible on host CPU. - * - * @Returns void - * - * @param void - * - * @warning __threadfence_system is a stub and map to no-op. - */ -//__device__ void __threadfence_system(void) __attribute__((deprecated("Provided with workaround -//configuration, see hip_kernel_language.md for details"))); -__device__ void __threadfence_system(void); - -// doxygen end Fence Fence -/** - * @} - */ - #endif // __HCC_OR_HIP_CLANG__ #if defined __HCC__ diff --git a/src/device_util.cpp b/src/device_util.cpp index 65ee5f4368..5107acd8c6 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -155,6 +155,3 @@ __host__ __device__ int max(int arg1, int arg2) { } __host__ void* __get_dynamicgroupbaseptr() { return nullptr; } - - -__device__ void __threadfence_system(void) { std::atomic_thread_fence(std::memory_order_seq_cst); } diff --git a/src/device_util.h b/src/device_util.h index 6603689d82..8fa96da9d9 100644 --- a/src/device_util.h +++ b/src/device_util.h @@ -125,7 +125,6 @@ __device__ double __hip_fast_dsqrt_rd(double x); __device__ double __hip_fast_dsqrt_rn(double x); __device__ double __hip_fast_dsqrt_ru(double x); __device__ double __hip_fast_dsqrt_rz(double x); -__device__ void __threadfence_system(void); float __hip_host_j0f(float x); double __hip_host_j0(double x); From 47d78e372e03e3c6d52081bee080f8712597ad82 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 5 Jul 2018 20:15:41 +0000 Subject: [PATCH 14/21] Implement min/max functions in HIP header Remove using hc::precise_math min and max. Instead we can use ocml directly for device and std:: for host. --- include/hip/hcc_detail/device_library_decls.h | 3 +++ include/hip/hcc_detail/hip_runtime.h | 11 +++++++++-- src/device_util.cpp | 7 ------- 3 files changed, 12 insertions(+), 9 deletions(-) diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 82c39b24f0..a7e81a1968 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -43,6 +43,9 @@ extern "C" __device__ float __ocml_rint_f32(float); extern "C" __device__ float __ocml_ceil_f32(float); extern "C" __device__ float __ocml_trunc_f32(float); +extern "C" __device__ float __ocml_fmin_f32(float, float); +extern "C" __device__ float __ocml_fmax_f32(float, float); + // Introduce local address space #define __local __attribute__((address_space(3))) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 3e0f2e27c5..48818fb15d 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -201,8 +201,15 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask #endif //__HIP_ARCH_GFX803__ == 1 -__host__ __device__ int min(int arg1, int arg2); -__host__ __device__ int max(int arg1, int arg2); +__device__ inline static int min(int arg1, int arg2) { + return (int)(__ocml_fmin_f32((float) arg1, (float) arg2)); +} +__device__ inline static int max(int arg1, int arg2) { + return (int)(__ocml_fmax_f32((float) arg1, (float) arg2)); +} + +__host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); } +__host__ inline static int max(int arg1, int arg2) { return std::max(arg1, arg2); } #endif // __HCC_OR_HIP_CLANG__ diff --git a/src/device_util.cpp b/src/device_util.cpp index 5107acd8c6..34e198c61d 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -147,11 +147,4 @@ __device__ void* __hip_hc_memset(void* dst, uint8_t val, size_t size) { // abort __device__ void abort() { return hc::abort(); } -__host__ __device__ int min(int arg1, int arg2) { - return (int)(hc::precise_math::fmin((float)arg1, (float)arg2)); -} -__host__ __device__ int max(int arg1, int arg2) { - return (int)(hc::precise_math::fmax((float)arg1, (float)arg2)); -} - __host__ void* __get_dynamicgroupbaseptr() { return nullptr; } From 76f86ef0971a624528346e90fab80384044ea755 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 5 Jul 2018 20:38:46 +0000 Subject: [PATCH 15/21] Implement hip_ldg Functions into HIP header Move all the function definitions for hip_ldg.cpp into hip_ldg.h header and enable for HIP clang path. --- CMakeLists.txt | 3 +- include/hip/hcc_detail/hip_ldg.h | 93 ++++++++++++++++++---------- src/hip_ldg.cpp | 83 ------------------------- tests/src/deviceLib/hip_test_ldg.cpp | 2 +- 4 files changed, 62 insertions(+), 119 deletions(-) delete mode 100644 src/hip_ldg.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 9095ff6531..a4da3b1920 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -193,8 +193,7 @@ if(HIP_PLATFORM STREQUAL "hcc") src/program_state.cpp) set(SOURCE_FILES_DEVICE - src/device_util.cpp - src/hip_ldg.cpp) + src/device_util.cpp) execute_process(COMMAND ${HCC_HOME}/bin/hcc-config --ldflags OUTPUT_VARIABLE HCC_LD_FLAGS) set(CMAKE_SHARED_LINKER_FLAGS "${CMAKE_SHARED_LINKER_FLAGS} ${HCC_LD_FLAGS} -Wl,-Bsymbolic") diff --git a/include/hip/hcc_detail/hip_ldg.h b/include/hip/hcc_detail/hip_ldg.h index 281550cd4c..a5b80b0a1b 100644 --- a/include/hip/hcc_detail/hip_ldg.h +++ b/include/hip/hcc_detail/hip_ldg.h @@ -23,54 +23,81 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H #define HIP_INCLUDE_HIP_HCC_DETAIL_HIP_LDG_H -#if defined __HCC__ -#if __hcc_workweek__ >= 16164 +#if defined(__HCC_OR_HIP_CLANG__) +#if __hcc_workweek__ >= 16164 || defined(__HIP_CLANG_ONLY__) #include "hip_vector_types.h" #include "host_defines.h" -__device__ char __ldg(const char*); -__device__ char2 __ldg(const char2*); -__device__ char4 __ldg(const char4*); -__device__ signed char __ldg(const signed char*); -__device__ unsigned char __ldg(const unsigned char*); +__device__ inline static char __ldg(const char* ptr) { return *ptr; } -__device__ short __ldg(const short*); -__device__ short2 __ldg(const short2*); -__device__ short4 __ldg(const short4*); -__device__ unsigned short __ldg(const unsigned short*); +__device__ inline static char2 __ldg(const char2* ptr) { return *ptr; } -__device__ int __ldg(const int*); -__device__ int2 __ldg(const int2*); -__device__ int4 __ldg(const int4*); -__device__ unsigned int __ldg(const unsigned int*); +__device__ inline static char4 __ldg(const char4* ptr) { return *ptr; } + +__device__ inline static signed char __ldg(const signed char* ptr) { return ptr[0]; } + +__device__ inline static unsigned char __ldg(const unsigned char* ptr) { return ptr[0]; } -__device__ long __ldg(const long*); -__device__ unsigned long __ldg(const unsigned long*); +__device__ inline static short __ldg(const short* ptr) { return ptr[0]; } -__device__ long long __ldg(const long long*); -__device__ longlong2 __ldg(const longlong2*); -__device__ unsigned long long __ldg(const unsigned long long*); +__device__ inline static short2 __ldg(const short2* ptr) { return ptr[0]; } -__device__ uchar2 __ldg(const uchar2*); -__device__ uchar4 __ldg(const uchar4*); +__device__ inline static short4 __ldg(const short4* ptr) { return ptr[0]; } -__device__ ushort2 __ldg(const ushort2*); +__device__ inline static unsigned short __ldg(const unsigned short* ptr) { return ptr[0]; } -__device__ uint2 __ldg(const uint2*); -__device__ uint4 __ldg(const uint4*); -__device__ ulonglong2 __ldg(const ulonglong2*); +__device__ inline static int __ldg(const int* ptr) { return ptr[0]; } -__device__ float __ldg(const float*); -__device__ float2 __ldg(const float2*); -__device__ float4 __ldg(const float4*); +__device__ inline static int2 __ldg(const int2* ptr) { return ptr[0]; } -__device__ double __ldg(const double*); -__device__ double2 __ldg(const double2*); +__device__ inline static int4 __ldg(const int4* ptr) { return ptr[0]; } -#endif // __hcc_workweek__ +__device__ inline static unsigned int __ldg(const unsigned int* ptr) { return ptr[0]; } -#endif // __HCC__ + +__device__ inline static long __ldg(const long* ptr) { return ptr[0]; } + +__device__ inline static unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; } + + +__device__ inline static long long __ldg(const long long* ptr) { return ptr[0]; } + +__device__ inline static longlong2 __ldg(const longlong2* ptr) { return ptr[0]; } + +__device__ inline static unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; } + + +__device__ inline static uchar2 __ldg(const uchar2* ptr) { return ptr[0]; } + +__device__ inline static uchar4 __ldg(const uchar4* ptr) { return ptr[0]; } + + +__device__ inline static ushort2 __ldg(const ushort2* ptr) { return ptr[0]; } + + +__device__ inline static uint2 __ldg(const uint2* ptr) { return ptr[0]; } + +__device__ inline static uint4 __ldg(const uint4* ptr) { return ptr[0]; } + + +__device__ inline static ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; } + + +__device__ inline static float __ldg(const float* ptr) { return ptr[0]; } + +__device__ inline static float2 __ldg(const float2* ptr) { return ptr[0]; } + +__device__ inline static float4 __ldg(const float4* ptr) { return ptr[0]; } + + +__device__ inline static double __ldg(const double* ptr) { return ptr[0]; } + +__device__ inline static double2 __ldg(const double2* ptr) { return ptr[0]; } + +#endif // __hcc_workweek__ || defined(__HIP_CLANG_ONLY__) + +#endif // defined(__HCC_OR_HIP_CLANG__) #endif // HIP_LDG_H diff --git a/src/hip_ldg.cpp b/src/hip_ldg.cpp deleted file mode 100644 index bf94c05571..0000000000 --- a/src/hip_ldg.cpp +++ /dev/null @@ -1,83 +0,0 @@ -/* -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. -*/ - -#include "hip/hcc_detail/hip_ldg.h" -#include "hip/hcc_detail/hip_vector_types.h" - -__device__ char __ldg(const char* ptr) { return *ptr; } - -__device__ char2 __ldg(const char2* ptr) { return *ptr; } - -__device__ char4 __ldg(const char4* ptr) { return *ptr; } - -__device__ signed char __ldg(const signed char* ptr) { return ptr[0]; } - -__device__ unsigned char __ldg(const unsigned char* ptr) { return ptr[0]; } - -__device__ short __ldg(const short* ptr) { return ptr[0]; } - -__device__ short2 __ldg(const short2* ptr) { return ptr[0]; } - -__device__ short4 __ldg(const short4* ptr) { return ptr[0]; } - -__device__ unsigned short __ldg(const unsigned short* ptr) { return ptr[0]; } - -__device__ int __ldg(const int* ptr) { return ptr[0]; } - -__device__ int2 __ldg(const int2* ptr) { return ptr[0]; } - -__device__ int4 __ldg(const int4* ptr) { return ptr[0]; } - -__device__ unsigned int __ldg(const unsigned int* ptr) { return ptr[0]; } - - -__device__ long __ldg(const long* ptr) { return ptr[0]; } - -__device__ unsigned long __ldg(const unsigned long* ptr) { return ptr[0]; } - -__device__ long long __ldg(const long long* ptr) { return ptr[0]; } - -__device__ longlong2 __ldg(const longlong2* ptr) { return ptr[0]; } - -__device__ unsigned long long __ldg(const unsigned long long* ptr) { return ptr[0]; } - -__device__ uchar2 __ldg(const uchar2* ptr) { return ptr[0]; } - -__device__ uchar4 __ldg(const uchar4* ptr) { return ptr[0]; } - -__device__ ushort2 __ldg(const ushort2* ptr) { return ptr[0]; } - -__device__ uint2 __ldg(const uint2* ptr) { return ptr[0]; } - -__device__ uint4 __ldg(const uint4* ptr) { return ptr[0]; } - -__device__ ulonglong2 __ldg(const ulonglong2* ptr) { return ptr[0]; } - -__device__ float __ldg(const float* ptr) { return ptr[0]; } - -__device__ float2 __ldg(const float2* ptr) { return ptr[0]; } - -__device__ float4 __ldg(const float4* ptr) { return ptr[0]; } - -__device__ double __ldg(const double* ptr) { return ptr[0]; } - -__device__ double2 __ldg(const double2* ptr) { return ptr[0]; } diff --git a/tests/src/deviceLib/hip_test_ldg.cpp b/tests/src/deviceLib/hip_test_ldg.cpp index 4dea81d9b9..7274baa92c 100644 --- a/tests/src/deviceLib/hip_test_ldg.cpp +++ b/tests/src/deviceLib/hip_test_ldg.cpp @@ -35,7 +35,7 @@ THE SOFTWARE. #include "hip/hip_vector_types.h" #include "test_common.h" -#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__) +#if (__hcc_workweek__ >= 16164) || defined(__HIP_PLATFORM_NVCC__) || defined(__HIP_CLANG_ONLY__) #define HIP_ASSERT(x) (assert((x) == hipSuccess)) From 22b60afa0d350ba7b3fe98ab4257a2949626c657 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Thu, 5 Jul 2018 20:49:47 +0000 Subject: [PATCH 16/21] Enable surface functions on HIP clang path Fix surface test on HIP clang path. --- include/hip/hcc_detail/hip_runtime.h | 2 +- include/hip/hcc_detail/surface_functions.h | 2 -- 2 files changed, 1 insertion(+), 3 deletions(-) diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 48818fb15d..267f970377 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -110,9 +110,9 @@ extern int HIP_TRACE_API; #include #include #include +#include #if __HCC__ #include -#include #endif // __HCC__ // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define. diff --git a/include/hip/hcc_detail/surface_functions.h b/include/hip/hcc_detail/surface_functions.h index 562cc440ed..607f221901 100644 --- a/include/hip/hcc_detail/surface_functions.h +++ b/include/hip/hcc_detail/surface_functions.h @@ -23,8 +23,6 @@ THE SOFTWARE. #ifndef HIP_INCLUDE_HIP_HCC_DETAIL_SURFACE_FUNCTIONS_H #define HIP_INCLUDE_HIP_HCC_DETAIL_SURFACE_FUNCTIONS_H -#include -#include #include #define __SURFACE_FUNCTIONS_DECL__ static __inline__ __device__ From e0390ddc9ffd1b933e2d52df5339ab00794cee6b Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Sat, 7 Jul 2018 23:30:39 -0400 Subject: [PATCH 17/21] Add declare of __get_dynamicgroupbaseptr for host compilation --- include/hip/hcc_detail/device_functions.h | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 716c51a887..0b7dd81d67 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -735,7 +735,6 @@ int64_t __lanemask_lt() } #ifdef __HIP_DEVICE_COMPILE__ - __device__ inline void* __get_dynamicgroupbaseptr() @@ -743,6 +742,10 @@ void* __get_dynamicgroupbaseptr() // Get group segment base pointer. return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); } +#else +__device__ +void* __get_dynamicgroupbaseptr(); +#endif // __HIP_DEVICE_COMPILE__ __device__ inline @@ -750,7 +753,6 @@ void *__amdgcn_get_dynamicgroupbaseptr() { return __get_dynamicgroupbaseptr(); } -#endif // __HIP_DEVICE_COMPILE__ // hip.amdgcn.bc - sync threads From 5ec7973a6f2e6f009087f5c319345f37aeb15a73 Mon Sep 17 00:00:00 2001 From: "Yaxun (Sam) Liu" Date: Mon, 9 Jul 2018 00:18:39 -0400 Subject: [PATCH 18/21] Add workaround __local_to_generic --- include/hip/hcc_detail/device_functions.h | 4 +++- include/hip/hcc_detail/device_library_decls.h | 2 -- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 0b7dd81d67..1938170ce4 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -734,13 +734,15 @@ int64_t __lanemask_lt() return ballot; } +__device__ inline void* __local_to_generic(void* p) { return p; } + #ifdef __HIP_DEVICE_COMPILE__ __device__ inline void* __get_dynamicgroupbaseptr() { // Get group segment base pointer. - return (char*)__local_to_generic(__to_local(__llvm_amdgcn_groupstaticsize())); + return (char*)__local_to_generic((void*)__to_local(__llvm_amdgcn_groupstaticsize())); } #else __device__ diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index a7e81a1968..53ad7595fe 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -53,8 +53,6 @@ extern "C" __device__ float __ocml_fmax_f32(float, float); __device__ inline static __local void* __to_local(unsigned x) { return (__local void*)x; } #endif //__HIP_DEVICE_COMPILE__ -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); From d1323f4f9b6ff118df5b75b7757dc89cb203cbd4 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 10 Jul 2018 17:56:57 +0000 Subject: [PATCH 19/21] Remove activelanemask asm using ockl and llvm instrinsics Replace implementation of __any and __all functions using OCKL functions and replaced __ballot implementation to use llvm intrinsic llvm.amdgcn.icmp.i32 instead of calls to __activelanemask_v4_b64_b1 which is not convergent. --- include/hip/hcc_detail/device_functions.h | 29 +++++-------------- include/hip/hcc_detail/device_library_decls.h | 2 ++ include/hip/hcc_detail/llvm_intrinsics.h | 2 ++ 3 files changed, 12 insertions(+), 21 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 1938170ce4..5466982878 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -662,51 +662,38 @@ void __named_sync(int a, int b) { __builtin_amdgcn_s_barrier(); } #endif // __HIP_DEVICE_COMPILE__ // warp vote function __all __any __ballot -extern "C" __device__ inline uint64_t __activelanemask_v4_b64_b1(unsigned int input) { - uint64_t output; - // define i64 @__activelanemask_v4_b64_b1(i32 %input) #5 { - // %a = tail call i64 asm "v_cmp_ne_i32_e64 $0, 0, $1", "=s,v"(i32 %input) #9 - // ret i64 %a - // } - __asm("v_cmp_ne_i32_e64 %0, 0, %1" : "=s"(output) : "v"(input)); - return output; -} - -__device__ -inline -unsigned int __activelanecount_u32_b1(unsigned int input) { - return __popcll(__activelanemask_v4_b64_b1(input)); -} - __device__ inline int __all(int predicate) { - return __popcll(__activelanemask_v4_b64_b1(predicate)) == __activelanecount_u32_b1(1); + return __ockl_wfall_i32(predicate); } __device__ inline int __any(int predicate) { #ifdef NVCC_COMPAT - if (__popcll(__activelanemask_v4_b64_b1(predicate)) != 0) + if (__ockl_wfany_i32(predicate) != 0) return 1; else return 0; #else - return __popcll(__activelanemask_v4_b64_b1(predicate)); + return __ockl_wfany_i32(predicate); #endif } +// XXX from llvm/include/llvm/IR/InstrTypes.h +#define ICMP_NE 33 + __device__ inline unsigned long long int __ballot(int predicate) { - return __activelanemask_v4_b64_b1(predicate); + return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE); } __device__ inline unsigned long long int __ballot64(int predicate) { - return __activelanemask_v4_b64_b1(predicate); + return __llvm_amdgcn_icmp_i32(predicate, 0, ICMP_NE); } // hip.amdgcn.bc - lanemask diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 53ad7595fe..64e4ff8898 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -30,6 +30,8 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" +extern "C" __device__ bool __ockl_wfany_i32(int); +extern "C" __device__ bool __ockl_wfall_i32(int); extern "C" __device__ int32_t __ockl_activelane_u32(void); extern "C" __device__ uint __ockl_mul24_u32(uint, uint); diff --git a/include/hip/hcc_detail/llvm_intrinsics.h b/include/hip/hcc_detail/llvm_intrinsics.h index 02df3c2fbe..6f2fc45626 100644 --- a/include/hip/hcc_detail/llvm_intrinsics.h +++ b/include/hip/hcc_detail/llvm_intrinsics.h @@ -31,6 +31,8 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" +__device__ ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32"); + __device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize"); From 1f871736e573f0a6ffec6ea83520e84e8630c9ad Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 10 Jul 2018 18:27:32 +0000 Subject: [PATCH 20/21] Add func attributes to match ocml and ockl --- include/hip/hcc_detail/device_library_decls.h | 28 +++++++++---------- 1 file changed, 14 insertions(+), 14 deletions(-) diff --git a/include/hip/hcc_detail/device_library_decls.h b/include/hip/hcc_detail/device_library_decls.h index 64e4ff8898..2a14b0b814 100644 --- a/include/hip/hcc_detail/device_library_decls.h +++ b/include/hip/hcc_detail/device_library_decls.h @@ -30,23 +30,23 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -extern "C" __device__ bool __ockl_wfany_i32(int); -extern "C" __device__ bool __ockl_wfall_i32(int); -extern "C" __device__ int32_t __ockl_activelane_u32(void); +extern "C" __device__ __attribute__((const)) bool __ockl_wfany_i32(int); +extern "C" __device__ __attribute__((const)) bool __ockl_wfall_i32(int); +extern "C" __device__ uint __ockl_activelane_u32(void); -extern "C" __device__ uint __ockl_mul24_u32(uint, uint); -extern "C" __device__ int __ockl_mul24_i32(int, int); -extern "C" __device__ uint __ockl_mul_hi_u32(uint, uint); -extern "C" __device__ int __ockl_mul_hi_i32(int, int); -extern "C" __device__ uint __ockl_sad_u32(uint, uint, uint); +extern "C" __device__ __attribute__((const)) uint __ockl_mul24_u32(uint, uint); +extern "C" __device__ __attribute__((const)) int __ockl_mul24_i32(int, int); +extern "C" __device__ __attribute__((const)) uint __ockl_mul_hi_u32(uint, uint); +extern "C" __device__ __attribute__((const)) int __ockl_mul_hi_i32(int, int); +extern "C" __device__ __attribute__((const)) uint __ockl_sad_u32(uint, uint, uint); -extern "C" __device__ float __ocml_floor_f32(float); -extern "C" __device__ float __ocml_rint_f32(float); -extern "C" __device__ float __ocml_ceil_f32(float); -extern "C" __device__ float __ocml_trunc_f32(float); +extern "C" __device__ __attribute__((const)) float __ocml_floor_f32(float); +extern "C" __device__ __attribute__((const)) float __ocml_rint_f32(float); +extern "C" __device__ __attribute__((const)) float __ocml_ceil_f32(float); +extern "C" __device__ __attribute__((const)) float __ocml_trunc_f32(float); -extern "C" __device__ float __ocml_fmin_f32(float, float); -extern "C" __device__ float __ocml_fmax_f32(float, float); +extern "C" __device__ __attribute__((const)) float __ocml_fmin_f32(float, float); +extern "C" __device__ __attribute__((const)) float __ocml_fmax_f32(float, float); // Introduce local address space #define __local __attribute__((address_space(3))) From 4a7ad93655954a81e1b9e863f9ae7fa2cc92e25d Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 10 Jul 2018 18:56:48 +0000 Subject: [PATCH 21/21] Fix min/max, icmp asm and add comment for conversion functions --- include/hip/hcc_detail/device_functions.h | 3 ++- include/hip/hcc_detail/hip_runtime.h | 4 ++-- include/hip/hcc_detail/llvm_intrinsics.h | 4 +++- 3 files changed, 7 insertions(+), 4 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 5466982878..47c25c87d7 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -440,7 +440,8 @@ __device__ static inline char4 __hip_hc_mul8pk(char4 in1, char4 in2) { } /* -Rounding modes are not yet supported in HIP + * Rounding modes are not yet supported in HIP + * TODO: Conversion functions are not correct, need to fix when BE is ready */ __device__ static inline float __double2float_rd(double x) { return (double)x; } diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index 267f970377..13eaf4fda4 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -202,10 +202,10 @@ __device__ int __hip_move_dpp(int src, int dpp_ctrl, int row_mask, int bank_mask #endif //__HIP_ARCH_GFX803__ == 1 __device__ inline static int min(int arg1, int arg2) { - return (int)(__ocml_fmin_f32((float) arg1, (float) arg2)); + return (arg1 < arg2) ? arg1 : arg2; } __device__ inline static int max(int arg1, int arg2) { - return (int)(__ocml_fmax_f32((float) arg1, (float) arg2)); + return (arg1 > arg2) ? arg1 : arg2; } __host__ inline static int min(int arg1, int arg2) { return std::min(arg1, arg2); } diff --git a/include/hip/hcc_detail/llvm_intrinsics.h b/include/hip/hcc_detail/llvm_intrinsics.h index 6f2fc45626..dc6fd05c52 100644 --- a/include/hip/hcc_detail/llvm_intrinsics.h +++ b/include/hip/hcc_detail/llvm_intrinsics.h @@ -31,7 +31,9 @@ THE SOFTWARE. #include "hip/hcc_detail/host_defines.h" -__device__ ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32"); +__device__ +__attribute__((convergent)) +ulong __llvm_amdgcn_icmp_i32(uint x, uint y, uint z) __asm("llvm.amdgcn.icmp.i32"); __device__ unsigned __llvm_amdgcn_groupstaticsize() __asm("llvm.amdgcn.groupstaticsize");