atomic min CAS, cuda implementation

atomic_min.cu

__device__ float customAtomicMin( unsigned long long int* address, unsigned long long int newVal )
{
    unsigned long long int old = *address, assumed;

    if( *((float*)&old ) <= *((float*)&newVal) ) 
        return *((float*)&old);
    do
    {
        assumed = old;
        if( *((float*)&assumed) <= *((float*)&newVal)  )
        {
            break;
        }
        old = atomicCAS( address, assumed, newVal );

    }
    while( assumed != old );

    return *((float*)&old);
}