Minor compare exchange optimization

Currently, `compare_exchange_strong` is using `__stronger_order_cuda`:
```cpp
inline __host__ __device__ int __stronger_order_cuda(int __a, int __b) {
int const __max = __a > __b ? __a : __b;
if(__max != __ATOMIC_RELEASE)
return __max;
static int const __xform[] = {
__ATOMIC_RELEASE,
__ATOMIC_ACQ_REL,
__ATOMIC_ACQ_REL,
__ATOMIC_RELEASE };
return __xform[__a < __b ? __a : __b];
}
```
The code above leads to actual memory loads. We can consider the following optimization:
```cpp
inline __host__ __device__ int __stronger_order_cuda(int __a, int __b) {
int const __max = __a > __b ? __a : __b;
if(__max != __ATOMIC_RELEASE)
return __max;
const int __min = __a < __b ? __a : __b;
return __ATOMIC_ACQ_REL - ((__min & 1) == (__min >> 1));
}
```
The change leads to about 4% better performance of compare exchange on mobile 3070 ti when memory ordering is not known at compile time:


When the memory ordering is known at compile time, there's no difference in generated SASS for both versions. Here's the benchmark:

```cpp
#include
#include

constexpr int threads_in_block = 1024;

__launch_bounds__(threads_in_block)
__global__ void kernel(int *ptr, int target, cuda::memory_order success, cuda::memory_order failure) {
__shared__ int cache;

int expected = -1;
if (threadIdx.x == target) {
cache = expected;
}
__syncthreads();

cuda::atomic_ref ref(cache);

if (ref.compare_exchange_strong(expected, threadIdx.x, success, failure)) {
ptr[blockIdx.x] = threadIdx.x;
}
}

int main() {
int blocks_in_grid = 256 * 1024;
int n = blocks_in_grid;

int *ptr{};
cudaMalloc(&ptr, sizeof(int) * n);
cudaMemset(ptr, 0, sizeof(int) * n);

cudaEvent_t begin, end;
cudaEventCreate(&begin);
cudaEventCreate(&end);

cudaEventRecord(begin);
kernel<<>>(ptr, 0, cuda::memory_order_release, cuda::memory_order_relaxed);
cudaEventRecord(end);
cudaEventSynchronize(end);

float ms{};
cudaEventElapsedTime(&ms, begin, end);

std::cout << ms << std::endl;

cudaEventDestroy(end);
cudaEventDestroy(begin);

cudaFree(ptr);
}
```
```bash
nvcc -gencode arch=compute_86,code=sm_86 -std=c++17 -DNDEBUG -O3 main.cu
```