diff --git a/hipamd/include/hip/amd_detail/amd_device_functions.h b/hipamd/include/hip/amd_detail/amd_device_functions.h index 44d4dc4dea..4fc5cde83b 100644 --- a/hipamd/include/hip/amd_detail/amd_device_functions.h +++ b/hipamd/include/hip/amd_detail/amd_device_functions.h @@ -118,25 +118,25 @@ __device__ static int32_t __fns64(uint64_t mask, uint32_t base, int32_t offset) return total; } -__device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset) { - uint64_t temp_mask = mask; +__device__ static int32_t __fns32(uint32_t mask, uint32_t base, int32_t offset) { + uint32_t temp_mask = mask; int32_t temp_offset = offset; if (offset == 0) { temp_mask &= (1 << base); temp_offset = 1; } else if (offset < 0) { - temp_mask = __builtin_bitreverse64(mask); - base = 63 - base; + temp_mask = __builtin_bitreverse32(mask); + base = 31 - base; temp_offset = -offset; } - temp_mask = temp_mask & ((~0ULL) << base); - if (__builtin_popcountll(temp_mask) < temp_offset) + temp_mask = temp_mask & ((~0U) << base); + if (__builtin_popcount(temp_mask) < temp_offset) return -1; int32_t total = 0; - for (int i = 0x20; i > 0; i >>= 1) { - uint64_t temp_mask_lo = temp_mask & ((1ULL << i) - 1); - int32_t pcnt = __builtin_popcountll(temp_mask_lo); + for (int i = 0x10; i > 0; i >>= 1) { + uint32_t temp_mask_lo = temp_mask & ((1U << i) - 1); + int32_t pcnt = __builtin_popcount(temp_mask_lo); if (pcnt < temp_offset) { temp_mask = temp_mask >> i; temp_offset -= pcnt; @@ -147,10 +147,16 @@ __device__ static int32_t __fns32(uint64_t mask, uint32_t base, int32_t offset) } } if (offset < 0) - return 63 - total; + return 31 - total; else return total; } + +// Wrapper around __fns32() to make porting from CUDA easier +__device__ static int32_t __fns(unsigned int mask, unsigned int base, int offset) { + return __fns32(mask, base, offset); +} + __device__ static inline unsigned int __brev(unsigned int input) { return __builtin_bitreverse32(input); }