Эффективная оценка функции отображения индексов в CUDA

Question

В ядре CUDA мне нужно найти ключ, который отображается в threadIdx.

отображение может выглядеть следующим образом:

ключ -> threadIdx

0 -> {0,1,2,3,4}

1 -> {5,6 , 7}

2 -> {8,9,10}

...

Каждый ключ k_i отображается на n_i (различный, произвольный n_i с n_i>0) нитками. Ключ будет использоваться для получения соответствующего значения в глобальном массиве. Это значение затем используется в последующих вычислениях в этом ядре.

Отображение может быть построена в виде кусочно-постоянной функции:

Количество клавиш не ограничено 3 (это только пример!), И только известны во время выполнения, а также соответствующей «ширине» каждой клавиши.

Как я могу эффективно узнать соответствующий ключ в ядре CUDA? Я подумал о двух следующих вариантов:

1. с использованием бинарного поиска внутри ядра (память эффективной)

2. предварительно рассчитав отображение для каждого threadIDx, то запуск ядра (выполнения эффективной)

   0 0 0 0 0 1 1 1 2 2 2 ...

Есть ли лучший способ достичь этого?

Answer 1

Существует еще один алгоритм, который дает вам что-то среднее с точки зрения производительности памяти и времени работы: Предположим, что общее количество потоков: N. Возьмем номер M, который близок к sqrt(N) и разделите все темы на группы на M нить каждой (последняя будет неполной). Теперь предварительно запрограммируйте ключ только для первого потока в каждой группе (их idxes будет 0, M, 2M и т. Д.). Это дает нам O(sqrt(N)) асимптотику памяти. Теперь в ядре мы можем легко найти индекс текущей группы (groupIdx = threadIdx/M) и следующей группы (groupIdx + 1). Для каждого из них мы знаем заранее рассчитанные ключи key[groupIdx] и key[groupIdx + 1]. Теперь вы можете сделать BS, но возьмите сегмент [key[groupIdx]; key[groupIdx + 1]] для поиска вместо [1; MAX_KEY_VALUES].

Answer 2

Вы можете комбинировать - каждый поток, когда он начинается, находит свой собственный ключ с BS и сохраняет его в массиве.