ff_storage.cuh

#pragma once
#include <cstdint>

using namespace std;

#define LIMBS_ALIGNMENT(x) ((x) % 4 == 0 ? 16 : ((x) % 2 == 0 ? 8 : 4))

template <unsigned LIMBS_COUNT> struct __align__(LIMBS_ALIGNMENT(LIMBS_COUNT)) ff_storage {
  static constexpr unsigned LC = LIMBS_COUNT;
  uint32_t limbs[LIMBS_COUNT];
};

template <unsigned LIMBS_COUNT> struct __align__(LIMBS_ALIGNMENT(LIMBS_COUNT)) ff_storage_wide {
  static_assert(LIMBS_COUNT ^ 1);
  static constexpr unsigned LC = LIMBS_COUNT;
  static constexpr unsigned LC2 = LIMBS_COUNT * 2;
  uint32_t limbs[LC2];

  void __device__ __forceinline__ set_lo(const ff_storage<LIMBS_COUNT> &in) {
#pragma unroll
    for (unsigned i = 0; i < LC; i++)
      limbs[i] = in.limbs[i];
  }

  void __device__ __forceinline__ set_hi(const ff_storage<LIMBS_COUNT> &in) {
#pragma unroll
    for (unsigned i = 0; i < LC; i++)
      limbs[i + LC].x = in.limbs[i];
  }

  ff_storage<LC> __device__ __forceinline__ get_lo() {
    ff_storage<LC> out{};
#pragma unroll
    for (unsigned i = 0; i < LC; i++)
      out.limbs[i] = limbs[i];
    return out;
  }

  ff_storage<LC> __device__ __forceinline__ get_hi() {
    ff_storage<LC> out{};
#pragma unroll
    for (unsigned i = 0; i < LC; i++)
      out.limbs[i] = limbs[i + LC].x;
    return out;
  }
};