From b445e450074d458f34e7ec048ca7e0504f3bcbb7 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 --- include/hip/hcc_detail/device_functions.h | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index 5e1c8a4e6b..7b147cc307 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/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; }