Spaces:
Running
Running
File size: 5,335 Bytes
c61ccee |
1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 82 83 84 85 86 87 88 89 90 91 92 93 94 95 96 97 98 99 100 101 102 103 104 105 106 107 108 109 110 111 112 113 114 115 116 117 118 119 120 121 122 |
#pragma once
#include <cuda.h>
#include <limits.h>
#include <math.h>
#include <float.h>
// NumericLimits.cuh is a holder for numeric limits definitions of commonly used
// types. This header is very specific to ROCm HIP and may be removed in the future.
// This header is derived from the legacy THCNumerics.cuh.
// The lower_bound and upper_bound constants are same as lowest and max for
// integral types, but are -inf and +inf for floating point types. They are
// useful in implementing min, max, etc.
namespace at {
template <typename T>
struct numeric_limits {
};
// WARNING: the following at::numeric_limits definitions are there only to support
// HIP compilation for the moment. Use std::numeric_limits if you are not
// compiling for ROCm.
// from @colesbury: "The functions on numeric_limits aren't marked with
// __device__ which is why they don't work with ROCm. CUDA allows them
// because they're constexpr."
namespace {
// ROCm doesn't like INFINITY too.
constexpr double inf = INFINITY;
}
template <>
struct numeric_limits<bool> {
static inline __host__ __device__ bool lowest() { return false; }
static inline __host__ __device__ bool max() { return true; }
static inline __host__ __device__ bool lower_bound() { return false; }
static inline __host__ __device__ bool upper_bound() { return true; }
};
template <>
struct numeric_limits<uint8_t> {
static inline __host__ __device__ uint8_t lowest() { return 0; }
static inline __host__ __device__ uint8_t max() { return UINT8_MAX; }
static inline __host__ __device__ uint8_t lower_bound() { return 0; }
static inline __host__ __device__ uint8_t upper_bound() { return UINT8_MAX; }
};
template <>
struct numeric_limits<int8_t> {
static inline __host__ __device__ int8_t lowest() { return INT8_MIN; }
static inline __host__ __device__ int8_t max() { return INT8_MAX; }
static inline __host__ __device__ int8_t lower_bound() { return INT8_MIN; }
static inline __host__ __device__ int8_t upper_bound() { return INT8_MAX; }
};
template <>
struct numeric_limits<int16_t> {
static inline __host__ __device__ int16_t lowest() { return INT16_MIN; }
static inline __host__ __device__ int16_t max() { return INT16_MAX; }
static inline __host__ __device__ int16_t lower_bound() { return INT16_MIN; }
static inline __host__ __device__ int16_t upper_bound() { return INT16_MAX; }
};
template <>
struct numeric_limits<int32_t> {
static inline __host__ __device__ int32_t lowest() { return INT32_MIN; }
static inline __host__ __device__ int32_t max() { return INT32_MAX; }
static inline __host__ __device__ int32_t lower_bound() { return INT32_MIN; }
static inline __host__ __device__ int32_t upper_bound() { return INT32_MAX; }
};
template <>
struct numeric_limits<int64_t> {
#ifdef _MSC_VER
static inline __host__ __device__ int64_t lowest() { return _I64_MIN; }
static inline __host__ __device__ int64_t max() { return _I64_MAX; }
static inline __host__ __device__ int64_t lower_bound() { return _I64_MIN; }
static inline __host__ __device__ int64_t upper_bound() { return _I64_MAX; }
#else
static inline __host__ __device__ int64_t lowest() { return INT64_MIN; }
static inline __host__ __device__ int64_t max() { return INT64_MAX; }
static inline __host__ __device__ int64_t lower_bound() { return INT64_MIN; }
static inline __host__ __device__ int64_t upper_bound() { return INT64_MAX; }
#endif
};
template <>
struct numeric_limits<at::Half> {
static inline __host__ __device__ at::Half lowest() { return at::Half(0xFBFF, at::Half::from_bits()); }
static inline __host__ __device__ at::Half max() { return at::Half(0x7BFF, at::Half::from_bits()); }
static inline __host__ __device__ at::Half lower_bound() { return at::Half(0xFC00, at::Half::from_bits()); }
static inline __host__ __device__ at::Half upper_bound() { return at::Half(0x7C00, at::Half::from_bits()); }
};
template <>
struct numeric_limits<at::BFloat16> {
static inline __host__ __device__ at::BFloat16 lowest() { return at::BFloat16(0xFF7F, at::BFloat16::from_bits()); }
static inline __host__ __device__ at::BFloat16 max() { return at::BFloat16(0x7F7F, at::BFloat16::from_bits()); }
static inline __host__ __device__ at::BFloat16 lower_bound() { return at::BFloat16(0xFF80, at::BFloat16::from_bits()); }
static inline __host__ __device__ at::BFloat16 upper_bound() { return at::BFloat16(0x7F80, at::BFloat16::from_bits()); }
};
template <>
struct numeric_limits<float> {
static inline __host__ __device__ float lowest() { return -FLT_MAX; }
static inline __host__ __device__ float max() { return FLT_MAX; }
static inline __host__ __device__ float lower_bound() { return -static_cast<float>(inf); }
static inline __host__ __device__ float upper_bound() { return static_cast<float>(inf); }
};
template <>
struct numeric_limits<double> {
static inline __host__ __device__ double lowest() { return -DBL_MAX; }
static inline __host__ __device__ double max() { return DBL_MAX; }
static inline __host__ __device__ double lower_bound() { return -inf; }
static inline __host__ __device__ double upper_bound() { return inf; }
};
} // namespace at
|