__inline__ __device__
T atomic_compare_exchange( volatile T * const dest , const T & compare ,
typename Kokkos::Impl::enable_if< sizeof(T) == sizeof(int) , const T & >::type val )
{
const int tmp = atomicCAS( (int*) dest , *((int*)&compare) , *((int*)&val) );
return *((T*)&tmp);
}
EPP_DEVICE void setSNSG(
float boxel_size, unsigned int capacity,
unsigned int x_boxel_count,
unsigned int y_boxel_count,
unsigned int z_boxel_count,
unsigned int *buffer,
const float3 &_position, unsigned int _index,
SliceIndex *slice_index
) {
int3 position_in_grid = emath::vector_cast< int3 >( _position / boxel_size );
position_in_grid.z = static_cast< unsigned int >( position_in_grid.z ) % z_boxel_count;
position_in_grid.x -= slice_index[ position_in_grid.z ].x_min;
position_in_grid.y -= slice_index[ position_in_grid.z ].y_min;
position_in_grid.x = static_cast< unsigned int >( position_in_grid.x ) % x_boxel_count;
position_in_grid.y = static_cast< unsigned int >( position_in_grid.y ) % y_boxel_count;
int where_to_set = (
slice_index[ position_in_grid.z ].offset +
position_in_grid.y * slice_index[ position_in_grid.z ].width +
position_in_grid.x ) * capacity;
int element_index;
_index++;
for( element_index = 0; element_index != capacity; element_index++ )
if( !atomicCAS( buffer + where_to_set + element_index, 0, _index ) )
break;
}
__inline__ __device__
T atomic_fetch_add( volatile T * const dest ,
typename Kokkos::Impl::enable_if< sizeof(T) == sizeof(int) , const T >::type val )
{
#ifdef KOKKOS_HAVE_CXX11
union U {
int i ;
T t ;
KOKKOS_INLINE_FUNCTION U() {};
} assume , oldval , newval ;
#else
union U {
int i ;
T t ;
} assume , oldval , newval ;
#endif
oldval.t = *dest ;
do {
assume.i = oldval.i ;
newval.t = assume.t + val ;
oldval.i = atomicCAS( (int*)dest , assume.i , newval.i );
} while ( assume.i != oldval.i );
return oldval.t ;
}
CUGIP_DECL_DEVICE inline float
atomicFloatCAS(float *address, float old, float val)
{
int i_val = __float_as_int(val);
int tmp0 = __float_as_int(old);
return __int_as_float(atomicCAS((int *)address, tmp0, i_val));
}
__inline__ __device__
T atomic_compare_exchange( volatile T * const dest , const T & compare ,
typename Kokkos::Impl::enable_if< sizeof(T) != sizeof(int) &&
sizeof(T) == sizeof(unsigned long long int) , const T & >::type val )
{
typedef unsigned long long int type ;
const type tmp = atomicCAS( (type*) dest , *((type*)&compare) , *((type*)&val) );
return *((T*)&tmp);
}
__device__ static double atomicMul(double* address, double val)
{
unsigned long long int* address_as_ull = (unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
old = atomicCAS(address_as_ull, assumed, __double_as_longlong(val * __longlong_as_double(assumed)));
} while (assumed != old);
return __longlong_as_double(old);
}
//float atomic func
__device__ static float atomicMul(float* address, float val)
{
int* address_as_int = (int*)address;
int old = *address_as_int, assumed;
do {
assumed = old;
old = atomicCAS(address_as_int, assumed, __float_as_int(__int_as_float(assumed) * val));
} while (assumed != old);
return __int_as_float(old);
}
__device__ static double atomicMin(double* address, double val)
{
unsigned long long int* address_as_ull = (unsigned long long int*)address;
unsigned long long int old = *address_as_ull, assumed;
do {
assumed = old;
if(__longlong_as_double(assumed) <= val) break;
old = atomicCAS(address_as_ull, assumed, val);
} while (assumed != old);
return __longlong_as_double(old);
}
__device__
bool atomically_obtain_ownership(int thread_idx)
{
const int no_owner = blockDim.x + 1;
// only grab the tag if it has no other owner
int old_owner = atomicCAS(m_owner, no_owner, thread_idx);
// we are the owner if it had no previous owner or if we already owned it
return (old_owner == no_owner) || (old_owner == thread_idx);
}
__inline__ __device__
T atomic_fetch_sub( volatile T * const dest ,
typename Kokkos::Impl::enable_if< sizeof(T) == sizeof(int) , const T >::type val )
{
union { int i ; T t ; } oldval , assume , newval ;
oldval.t = *dest ;
do {
assume.i = oldval.i ;
newval.t = assume.t - val ;
oldval.i = atomicCAS( (int*)dest , assume.i , newval.i );
} while ( assumed.i != oldval.i );
return oldval.t ;
}
int moveWatch (watchList **w, watchList *t) {
watchList *local;
t->tail = NULL;
do {
local = atomicRead(*w);
if (local == doNotAdd) {
return 0;
} else {
t->tail = local;
}
} while (!atomicCAS(*w,local,t)); /* OPT : Delay loop */
return 1;
}
__global__ void HIP_FUNCTION(testKernel,int *g_odata)
{
// access thread id
const unsigned int tid = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
// Test various atomic instructions
// Arithmetic atomic instructions
// Atomic addition
atomicAdd(&g_odata[0], 10);
// Atomic subtraction (final should be 0)
atomicSub(&g_odata[1], 10);
// Atomic exchange
atomicExch(&g_odata[2], tid);
// Atomic maximum
atomicMax(&g_odata[3], tid);
// Atomic minimum
atomicMin(&g_odata[4], tid);
// Atomic increment (modulo 17+1)
//atomicInc((unsigned int *)&g_odata[5], 17);
atomicInc((unsigned int *)&g_odata[5]);
// Atomic decrement
// atomicDec((unsigned int *)&g_odata[6], 137);
atomicDec((unsigned int *)&g_odata[6]);
// Atomic compare-and-swap
atomicCAS(&g_odata[7], tid-1, tid);
// Bitwise atomic instructions
// Atomic AND
atomicAnd(&g_odata[8], 2*tid+7);
// Atomic OR
atomicOr(&g_odata[9], 1 << tid);
// Atomic XOR
atomicXor(&g_odata[10], tid);
}
inline
__device__
void cuda_internal_atomic_add( double & update , double input )
{
typedef unsigned long long int UInt64 ;
UInt64 * const address = reinterpret_cast<UInt64*>( & update );
UInt64 test ;
union UType { double d ; UInt64 i ; } value ;
value.i = *address ; // Read existing value
do {
test = value.i ;
value.d += input ;
value.i = atomicCAS( address , test , value.i );
} while ( value.i != test );
}
KOKKOS_INLINE_FUNCTION
typename Kokkos::Impl::UnionPair<T,int,unsigned long long int>::first_type
atomic_fetch_add( volatile T * const dest , const T val )
{
typedef Kokkos::Impl::UnionPair<T,int,unsigned long long int> union_type ;
typedef typename union_type::second_type type ;
union_type assumed , old , newval ;
old.first = *dest ;
do {
assumed.second = old.second ;
newval.first = assumed.first + val ;
old.second = atomicCAS( (type *) union_type::cast( dest ),
assumed.second ,
newval.second );
} while ( assumed.second != old.second );
return old.first ;
}
__inline__ __device__
T atomic_fetch_add( volatile T * const dest ,
typename Kokkos::Impl::enable_if< sizeof(T) != sizeof(int) &&
sizeof(T) == sizeof(unsigned long long int) , const T >::type val )
{
union U {
unsigned long long int i ;
T t ;
KOKKOS_INLINE_FUNCTION U() {};
} assume , oldval , newval ;
oldval.t = *dest ;
do {
assume.i = oldval.i ;
newval.t = assume.t + val ;
oldval.i = atomicCAS( (unsigned long long int*)dest , assume.i , newval.i );
} while ( assume.i != oldval.i );
return oldval.t ;
}
__inline__ __device__
unsigned long long int atomic_compare_exchange( volatile unsigned long long int * const dest ,
const unsigned long long int compare ,
const unsigned long long int val )
{ return atomicCAS((unsigned long long int*)dest,compare,val); }
__inline__ __device__
int atomic_compare_exchange( volatile int * const dest, const int compare, const int val)
{ return atomicCAS((int*)dest,compare,val); }
