message: Manually define `GpuAtomic{Max,Min}` for `long long`.
  CUDA only defines `atomic{Max,Min}` for `long long` for compute capability 3.5 or higher. Hence, to prevent compile errors due to failed instantiation when compiling on `sm30` (or other similar ones), we manually define the corresponding wrappers.
diff: @@ -820,6 +820,11 @@ __device__ inline tensorflow::uint64 GpuAtomicMax(tensorflow::uint64* ptr,
  return detail::GpuAtomicCasHelper(
  ptr, [value](tensorflow::uint64 a) { return max(a, value); });
  }
  +
  +__device__ inline int64 GpuAtomicMax(int64* ptr, int64 value) {
  + return detail::GpuAtomicCasHelper(ptr,
  + [value](int64 a) { return max(a, value); });
  +}
  #endif
  CREATE_CUDA_DEVICE_FUNCTION_ALIAS(GpuAtomicMax, CudaAtomicMax);
  @@ -885,6 +890,11 @@ __device__ inline tensorflow::uint64 GpuAtomicMin(tensorflow::uint64* ptr,
  return detail::GpuAtomicCasHelper(
  ptr, [value](tensorflow::uint64 a) { return min(a, value); });
  }
  +
  +__device__ inline int64 GpuAtomicMin(int64* ptr, int64 value) {
  + return detail::GpuAtomicCasHelper(ptr,
  + [value](int64 a) { return min(a, value); });
  +}
  #endif
  CREATE_CUDA_DEVICE_FUNCTION_ALIAS(GpuAtomicMin, CudaAtomicMin);
  