AI/rocm-systems
2
0
Fork 0
Código Incidencias Pull Requests Acciones Paquetes Proyectos Lanzamientos Wiki Actividad

v1: Working on Integer Intrinsics

Explorar el Código 
1. Half way through
2. May not work
3. No test written

Change-Id: I705b743a78b142ff068e2521870e73fca7ad2b1c


[ROCm/hip commit: b09ad764a1]
Este commit está contenido en:
Aditya Atluri
2017-01-16 14:55:29 -06:00
padre 38edad98a6
commit 72202cccbd
Se han modificado 5 ficheros con 242 adiciones y 134 borrados
Descargar archivo de parche Descargar archivo de diferencias Expandir todos los archivos Contraer todos los archivos
projects/hip/src/device_util.cpp
-110
Ver fichero
@@ -1843,117 +1843,7 @@ __device__ unsigned int atomicDec(unsigned int* address,
return hc::__atomic_wrapdec(address,val);
}
//__mul24 __umul24
__device__ int __mul24(int arg1,
int arg2)
{
return hc::__mul24(arg1, arg2);
}
__device__ unsigned int __umul24(unsigned int arg1,
unsigned int arg2)
{
return hc::__mul24(arg1, arg2);
}
__device__ unsigned int test__popc(unsigned int input)
{
return hc::__popcount_u32_b32(input);
}
// integer intrinsic function __poc __clz __ffs __brev
__device__ unsigned int __popc( unsigned int input)
{
return hc::__popcount_u32_b32(input);
}
__device__ unsigned int test__popc(unsigned int input);
__device__ unsigned int __popcll( unsigned long long int input)
{
return hc::__popcount_u32_b64(input);
}
__device__ unsigned int __clz(unsigned int input)
{
#ifdef NVCC_COMPAT
return input == 0 ? 32 : hc::__firstbit_u32_u32( input);
#else
return hc::__firstbit_u32_u32( input);
#endif
}
__device__ unsigned int __clzll(unsigned long long int input)
{
#ifdef NVCC_COMPAT
return input == 0 ? 64 : hc::__firstbit_u32_u64( input);
#else
return hc::__firstbit_u32_u64( input);
#endif
}
__device__ unsigned int __clz( int input)
{
#ifdef NVCC_COMPAT
return input == 0 ? 32 : hc::__firstbit_u32_s32( input);
#else
return hc::__firstbit_u32_s32( input);
#endif
}
__device__ unsigned int __clzll( long long int input)
{
#ifdef NVCC_COMPAT
return input == 0 ? 64 : hc::__firstbit_u32_s64( input);
#else
return hc::__firstbit_u32_s64( input);
#endif
}
__device__ unsigned int __ffs(unsigned int input)
{
#ifdef NVCC_COMPAT
return hc::__lastbit_u32_u32( input)+1;
#else
return hc::__lastbit_u32_u32( input);
#endif
}
__device__ unsigned int __ffsll(unsigned long long int input)
{
#ifdef NVCC_COMPAT
return hc::__lastbit_u32_u64( input)+1;
#else
return hc::__lastbit_u32_u64( input);
#endif
}
__device__ unsigned int __ffs( int input)
{
#ifdef NVCC_COMPAT
return hc::__lastbit_u32_s32( input)+1;
#else
return hc::__lastbit_u32_s32( input);
#endif
}
__device__ unsigned int __ffsll( long long int input)
{
#ifdef NVCC_COMPAT
return hc::__lastbit_u32_s64( input)+1;
#else
return hc::__lastbit_u32_s64( input);
#endif
}
__device__ unsigned int __brev( unsigned int input)
{
return hc::__bitrev_b32( input);
}
__device__ unsigned long long int __brevll( unsigned long long int input)
{
return hc::__bitrev_b64( input);
}
// warp vote function __all __any __ballot
__device__ int __all( int input)