скомпилировать постоянный массив памяти в немедленное значение в CUDA

, пожалуйста, уважайте работу и добавили немного CUDA.

А, вот и ты:

Вы хотите, чтобы компилятор делал всю грязную работу за вас, и поэтому вы должны все и вся оценивать во время компиляции:

Сначала нам нужен статический массив, в котором мы можем хранить коэффициенты:

template <unsigned int index, long long... remPack> struct getValue;

template <unsigned int index, long long In, long long... remPack>
struct getValue<index, In, remPack...> {
  static const long long value = getValue<index - 1, remPack...>::value;
};

template <long long In, long long... remPack>
struct getValue<1, In, remPack...> {
  static const long long value = In;
};

template <long long... T> struct static_array {
  template <int idx> static __device__ int get() { return getValue<idx, T...>::value; }
};

Этот static_array хранит значения в системе типов C++ как long long. Я вернусь к этому позже в ответе.

Далее следует развернуть цикл for. Опять же, для этого используйте метапрограммирование шаблонов:

template <int First, int Last, template <int> class Fn> struct static_for {
  __device__ double operator()(double x, double y) const {
      return static_for<First + 1, Last, Fn>()(x, Fn<First + 1>()(x, y));
  }
};

template <int N, template <int> class Fn> struct static_for<N, N, Fn> {
  __device__ double operator()(double x, double y) const { return y; }
};

Поскольку мы делаем все статические во время компиляции, необходимо перемещать вход и выход каждого «цикла» через параметры и возвращаемое выражение operator().

Это решение очень статично, и с помощью большего количества метапрограмм шаблонов вы можете улучшить его.

Хорошо, теперь самое интересное. Расчеты:

template <int i> struct Function {
  __device__ double operator()(double x, double y) {
    double c = __longlong_as_double(static_array<12, 34, 22, 55, 24>::get<i>());
    return y * x + c;
  }
};

__device__ double some_function(double x) {
  return static_for<0, 5, Function>()(x, 0.0);
}

Система типов C++ допускает только целые типы в качестве нетиповых параметров шаблона, поэтому мы должны хранить наши doubles в long longs, а затем использовать функцию CUDA __longlong_as_double(), чтобы преобразовать их обратно. Это то, что мы должны принять на данный момент, и это может стать для вас препятствием, потому что это уже не «просто». Однако написать преобразователь double в long long не так уж сложно.

Все вычисления заключены в функторном объекте, который получает счетчик поездок из нашего static_loop в качестве аргумента шаблона. С помощью этого «счетчика пути» времени компиляции мы можем получить доступ к static_array преобразованию long long версии двойного возврата и вычислить FMA.

Благодаря компилятору CUDA (который здесь действительно хорошо справляется), это код PTX (nvcc -ptx -arch=sm_35 test.cu). Я использовал версию 7.0 RC1:

.visible .func  (.param .b64 func_retval0) _Z13some_functiond(
        .param .b64 _Z13some_functiond_param_0
)
{
        .reg .f64       %fd<7>;

        ld.param.f64    %fd1, [_Z13some_functiond_param_0];
        fma.rn.f64      %fd2, %fd1, 0d0000000000000000, 0d000000000000000C;
        fma.rn.f64      %fd3, %fd2, %fd1, 0d0000000000000022;
        fma.rn.f64      %fd4, %fd3, %fd1, 0d0000000000000016;
        fma.rn.f64      %fd5, %fd4, %fd1, 0d0000000000000037;
        fma.rn.f64      %fd6, %fd5, %fd1, 0d0000000000000018;
        st.param.f64    [func_retval0+0], %fd6;
        ret;
}