11#ifndef EIGEN_TENSOR_TENSOR_INTDIV_H
12#define EIGEN_TENSOR_TENSOR_INTDIV_H
15#include "./InternalHeaderCheck.h"
23EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<
sizeof(T) == 4,
int> count_leading_zeros(
const T val) {
24#ifdef EIGEN_GPU_COMPILE_PHASE
26#elif defined(SYCL_DEVICE_ONLY)
27 return cl::sycl::clz(val);
30 _BitScanReverse(&index, val);
33 EIGEN_STATIC_ASSERT(
sizeof(
unsigned long long) == 8, YOU_MADE_A_PROGRAMMING_MISTAKE);
34 return __builtin_clz(
static_cast<uint32_t
>(val));
39EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<
sizeof(T) == 8,
int> count_leading_zeros(
const T val) {
40#ifdef EIGEN_GPU_COMPILE_PHASE
42#elif defined(SYCL_DEVICE_ONLY)
43 return static_cast<int>(cl::sycl::clz(val));
44#elif EIGEN_COMP_MSVC && EIGEN_ARCH_x86_64
46 _BitScanReverse64(&index, val);
50 unsigned int lo = (
unsigned int)(val & 0xffffffff);
51 unsigned int hi = (
unsigned int)((val >> 32) & 0xffffffff);
54 n = 32 + count_leading_zeros<unsigned int>(lo);
56 n = count_leading_zeros<unsigned int>(hi);
59 EIGEN_STATIC_ASSERT(
sizeof(
unsigned long long) == 8, YOU_MADE_A_PROGRAMMING_MISTAKE);
60 return __builtin_clzll(
static_cast<uint64_t
>(val));
65struct UnsignedTraits {
66 typedef std::conditional_t<
sizeof(T) == 8, uint64_t, uint32_t> type;
71 typedef typename UnsignedTraits<T>::type type;
72 static constexpr int N =
sizeof(T) * 8;
76EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint32_t muluh(
const uint32_t a,
const T b) {
77#if defined(EIGEN_GPU_COMPILE_PHASE)
78 return __umulhi(a, b);
79#elif defined(SYCL_DEVICE_ONLY)
80 return cl::sycl::mul_hi(a,
static_cast<uint32_t
>(b));
82 return (
static_cast<uint64_t
>(a) * b) >> 32;
87EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint64_t muluh(
const uint64_t a,
const T b) {
88#if defined(EIGEN_GPU_COMPILE_PHASE)
89 return __umul64hi(a, b);
90#elif defined(SYCL_DEVICE_ONLY)
91 return cl::sycl::mul_hi(a,
static_cast<uint64_t
>(b));
92#elif EIGEN_COMP_MSVC && (EIGEN_ARCH_x86_64 || EIGEN_ARCH_ARM64)
93 return __umulh(a,
static_cast<uint64_t
>(b));
94#elif EIGEN_HAS_BUILTIN_INT128
95 __uint128_t v =
static_cast<__uint128_t
>(a) *
static_cast<__uint128_t
>(b);
96 return static_cast<uint64_t
>(v >> 64);
98 return (TensorUInt128<static_val<0>, uint64_t>(a) * TensorUInt128<static_val<0>, uint64_t>(b)).upper();
102template <
int N,
typename T>
103struct DividerHelper {
104 static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint32_t computeMultiplier(
const int log_div,
const T divider) {
105 EIGEN_STATIC_ASSERT(N == 32, YOU_MADE_A_PROGRAMMING_MISTAKE);
106 return static_cast<uint32_t
>((
static_cast<uint64_t
>(1) << (N + log_div)) / divider -
107 (
static_cast<uint64_t
>(1) << N) + 1);
112struct DividerHelper<64, T> {
113 static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE uint64_t computeMultiplier(
const int log_div,
const T divider) {
114#if EIGEN_HAS_BUILTIN_INT128 && !defined(EIGEN_GPU_COMPILE_PHASE) && !defined(SYCL_DEVICE_ONLY)
115 return static_cast<uint64_t
>((
static_cast<__uint128_t
>(1) << (64 + log_div)) /
static_cast<__uint128_t
>(divider) -
116 (
static_cast<__uint128_t
>(1) << 64) + 1);
118 const uint64_t shift = 1ULL << log_div;
119 TensorUInt128<uint64_t, uint64_t> result =
120 TensorUInt128<uint64_t, static_val<0> >(shift, 0) / TensorUInt128<static_val<0>, uint64_t>(divider) -
121 TensorUInt128<static_val<1>, static_val<0> >(1, 0) + TensorUInt128<static_val<0>, static_val<1> >(1);
122 return static_cast<uint64_t
>(result);
138template <
typename T,
bool div_gt_one = false>
139struct TensorIntDivisor {
141 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorIntDivisor() =
default;
146 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorIntDivisor(
const T divider) {
147 const int N = DividerTraits<T>::N;
148 eigen_assert(
static_cast<typename UnsignedTraits<T>::type
>(divider) < NumTraits<UnsignedType>::highest() / 2);
149 eigen_assert(divider > 0);
152 const int leading_zeros = count_leading_zeros(
static_cast<UnsignedType
>(divider));
153 int log_div = N - leading_zeros;
155 if ((
static_cast<typename UnsignedTraits<T>::type
>(1) << (log_div - 1)) ==
156 static_cast<typename UnsignedTraits<T>::type
>(divider))
159 multiplier = DividerHelper<N, T>::computeMultiplier(log_div, divider);
160 shift1 = log_div > 1 ? 1 : log_div;
161 shift2 = log_div > 1 ? log_div - 1 : 0;
166 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T divide(
const T numerator)
const {
167 eigen_assert(
static_cast<typename UnsignedTraits<T>::type
>(numerator) < NumTraits<UnsignedType>::highest() / 2);
169 UnsignedType t1 = muluh(multiplier, numerator);
170 UnsignedType t = (
static_cast<UnsignedType
>(numerator) - t1) >> shift1;
171 return (t1 + t) >> shift2;
175 typedef typename DividerTraits<T>::type UnsignedType;
176 UnsignedType multiplier = 0;
185class TensorIntDivisor<int32_t, true> {
187 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE TensorIntDivisor() =
default;
189 EIGEN_DEVICE_FUNC TensorIntDivisor(int32_t divider) {
190 eigen_assert(divider >= 2);
194 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
int divide(
const int32_t n)
const {
195#ifdef EIGEN_GPU_COMPILE_PHASE
196 return __umulhi(magic, n) >> shift;
197#elif defined(SYCL_DEVICE_ONLY)
198 return cl::sycl::mul_hi(magic,
static_cast<uint32_t
>(n)) >> shift;
200 uint64_t v =
static_cast<uint64_t
>(magic) *
static_cast<uint64_t
>(n);
201 return static_cast<uint32_t
>(v >> 32) >> shift;
208 EIGEN_DEVICE_FUNC
void calcMagic(int32_t d) {
209 const unsigned two31 = 0x80000000;
211 unsigned t = two31 + (ad >> 31);
212 unsigned anc = t - 1 - t % ad;
214 unsigned q1 = two31 / anc;
215 unsigned r1 = two31 - q1 * anc;
216 unsigned q2 = two31 / ad;
217 unsigned r2 = two31 - q2 * ad;
234 }
while (q1 < delta || (q1 == delta && r1 == 0));
236 magic = (unsigned)(q2 + 1);
244template <
typename T,
bool div_gt_one>
245EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE T operator/(
const T& numerator,
const TensorIntDivisor<T, div_gt_one>& divisor) {
246 return divisor.divide(numerator);
Namespace containing all symbols from the Eigen library.