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