Альтернатива динамическому распределению небольших массивов внутри ядра?

Question

У меня есть ядро, для работы которого требуется небольшой объем памяти. Тем не менее, эта память зависит от ввода, поэтому я динамически выделяю ее с malloc внутри, но часто это очень маленький вектор, редко больше 5, почти никогда не больше 10. Однако каждый раз, когда я выполняю программу существует только 1 значение для длины, оно должно быть только динамическим для поддержки различных входных данных, длина одинакова при каждом выполнении всей программы.

(Короче говоря, это поиск в глубине первого дерева, где искомый объект может быть частью нескольких листьев, поэтому мне нужно отслеживать, какой из них является текущим путем поиска и сколько я проверил в каждом узле пути. Не актуально, просто добавив немного информации).

Я проверил разницу в скорости между динамическим выделением и статическим заданием размера (int path[6];), и неудивительно, что статическая версия примерно в 15 раз быстрее.

Есть ли альтернатива динамическому распределению массивов для моего конкретного случая (я знаю, что он всегда одинаковой длины для данного выполнения программы)? Я думал просто статически выделить вектор, например, 10 значений, так что я всегда в безопасности, но это накладывает произвольный верхний предел на мой размер ввода, который, даже если его трудно достичь, я бы предпочел не делать. Есть идеи?

Answer 1

Учитывая, что

1. Вы знаете значение требуемого размера локального хранилища a priori и
2. Вы знаете, что диапазон возможных локальных размеров хранилища мал и ограничен,

логическим решением является использование шаблонов с размером, переданным в качестве параметра шаблона:

template <int lrrysz>
__global__ void thekernel(float *in, float *out)
{

    float local[lrrysz];

    // your code goes here

}

template __global__ void thekernel<5>(float*, float*);
template __global__ void thekernel<6>(float*, float*);
template __global__ void thekernel<7>(float*, float*);
template __global__ void thekernel<8>(float*, float*);
template __global__ void thekernel<9>(float*, float*);
template __global__ void thekernel<10>(float*, float*);

Это дает другие преимущества, в частности, то, что размер массива теперь является постоянной времени компиляции. Это означает, что компилятор может развернуть множество оптимизаций, например развернуть циклы, и, если вам очень повезет, выложить массив в регистры. Это также означает, что обычные элементы управления компилятора, такие как границы запуска, пределы регистров и т. Д., Также работают должным образом и могут предлагать улучшения производительности при каждом различном размере локального хранилища.