From 0fca41a7ce770472dc2ad132ac3040f12a6014e2 Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Tue, 19 Jun 2018 19:22:31 +0000 Subject: [PATCH] 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. [ROCm/hip commit: e02fc7e680059047ab6bd7db29c39de017e0cf84] --- projects/hip/CMakeLists.txt | 3 +- .../include/hip/hcc_detail/device_functions.h | 475 ++++++++++++++---- .../hip/hcc_detail/device_library_decls.h | 5 + .../include/hip/hcc_detail/llvm_intrinsics.h | 29 ++ projects/hip/src/device_functions.cpp | 387 -------------- 5 files changed, 422 insertions(+), 477 deletions(-) delete mode 100644 projects/hip/src/device_functions.cpp diff --git a/projects/hip/CMakeLists.txt b/projects/hip/CMakeLists.txt index 537764548f..b649b20c21 100644 --- a/projects/hip/CMakeLists.txt +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/device_functions.h b/projects/hip/include/hip/hcc_detail/device_functions.h index 8bae5325fd..b147cd9b80 100644 --- a/projects/hip/include/hip/hcc_detail/device_functions.h +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/device_library_decls.h b/projects/hip/include/hip/hcc_detail/device_library_decls.h index b3fa556bd4..8bf3ce3a19 100644 --- a/projects/hip/include/hip/hcc_detail/device_library_decls.h +++ b/projects/hip/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/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h b/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h index 2c7819b535..02df3c2fbe 100644 --- a/projects/hip/include/hip/hcc_detail/llvm_intrinsics.h +++ b/projects/hip/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/projects/hip/src/device_functions.cpp b/projects/hip/src/device_functions.cpp deleted file mode 100644 index 8ef19bab3f..0000000000 --- a/projects/hip/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; -}