From 13f4ace668798da5fa35195e4c08a946c62fdbcd Mon Sep 17 00:00:00 2001 From: Aaron Enye Shi Date: Mon, 20 Aug 2018 16:46:12 +0000 Subject: [PATCH] Add bitextract and bitinsert functions [ROCm/clr commit: fb1da42d0e28f3512e73c17b69b38f91435cd78a] --- .../include/hip/hcc_detail/device_functions.h | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h index 5e1c8a4e6b..7b147cc307 100644 --- a/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h +++ b/projects/clr/hipamd/include/hip/hcc_detail/device_functions.h @@ -63,6 +63,22 @@ __device__ static inline unsigned int __clzll(long long int input) { return input == 0 ? 64 : input > 0 ? __builtin_clzl(input) : __builtin_clzl(~input); } +__device__ static inline unsigned int __lastbit_u32_u64(unsigned long long int input) { + return input == 0 ? -1 : __builtin_ctzl(input); +} + +__device__ static inline unsigned int __bitextract_u32(unsigned int src0, unsigned int src1, unsigned int src2) { + return (src0 << (32 - src1 - src2)) >> (32 - src2); +} + +__device__ static inline uint64_t __bitextract_u64(uint64_t src0, unsigned int src1, unsigned int src2) { + return (src0 << (64 - src1 - src2)) >> (64 - src2); +} + +__device__ static inline unsigned int __bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int src2, unsigned int src3) { + return (src0 & ~(((1 << src3) - 1) << src2)) | ((src1 & ((1 << src3) - 1)) << src2); +} + __device__ static inline unsigned int __ffs(unsigned int input) { return ( input == 0 ? -1 : __builtin_ctz(input) ) + 1; }