AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Clean up and add meaningful variables to bit funcs

Telusuri Sumber 
[ROCm/hip commit: b939a3c6ed]
This commit is contained in:
Aaron Enye Shi
2018-08-28 18:19:48 +00:00
orang tua 1d29819850
melakukan eb2e7e0671
1 mengubah file dengan 26 tambahan dan 17 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
projects/hip/include/hip/hcc_detail/device_functions.h
+26 -17
Melihat File
@@ -63,22 +63,6 @@ __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;
}
@@ -95,12 +79,37 @@ __device__ static inline unsigned int __ffsll(long long int input) {
return ( input == 0 ? -1 : __builtin_ctzl(input) ) + 1;
}
__device__ static inline unsigned int __brev(unsigned int input) { return __llvm_bitrev_b32(input); }
__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 inline unsigned int __lastbit_u32_u64(uint64_t input) {
return input == 0 ? -1 : __builtin_ctzl(input);
}
__device__ static inline unsigned int __bitextract_u32(unsigned int src0, unsigned int src1, unsigned int src2) {
uint32_t offset = src1 & 31;
uint32_t width = src2 & 31;
return width == 0 ? 0 : (src0 << (32 - offset - width)) >> (32 - width);
}
__device__ static inline uint64_t __bitextract_u64(uint64_t src0, unsigned int src1, unsigned int src2) {
uint64_t offset = src1 & 63;
uint64_t width = src2 & 63;
return width == 0 ? 0 : (src0 << (64 - offset - width)) >> (64 - width);
}
__device__ static inline unsigned int __bitinsert_u32(unsigned int src0, unsigned int src1, unsigned int src2, unsigned int src3) {
uint32_t offset = src2 & 31;
uint32_t width = src3 & 31;
uint32_t mask = (1 << width) - 1;
return ((src0 & ~(mask << offset)) | ((src1 & mask) << offset));
}
__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);