Для текущего проекта OpenCL GPGPU мне нужно отсортировать элементы в массиве по некоторому ключу с 64 возможными значениями. Мне нужно, чтобы в последнем массиве все элементы с одним и тем же ключом были смежными. Достаточно иметь ассоциативный массив new_index[old_index] в качестве выхода этой задачи.
Я разделил задачу на две части. Сначала я подсчитываю для каждого возможного ключа (корзины) количество элементов с этим ключом (которые попадают в эту корзину). Я просматриваю этот массив (генерирую сумму префиксов), который указывает новый диапазон индексов элементов для каждой корзины, например "начальные" индексы для каждой корзины.
На втором этапе каждому элементу необходимо присвоить новый индекс. Если бы я реализовал это на процессоре, алгоритм был бы примерно таким:
for all elements e:
new_index[e] = bucket_start[bucket(e)]++
Конечно, на GPU это не работает. Каждый элемент должен получить доступ к массиву bucket_start в режиме чтения-записи, который, по сути, является синхронизацией между всеми рабочими элементами, худшее, что мы можем сделать.
Идея состоит в том, чтобы провести вычисления в рабочих группах. Но я не уверен, как именно это должно быть сделано, поскольку у меня нет опыта в вычислениях GPGPU.
В глобальной памяти у нас есть массив начала сегмента, инициализированный префиксной суммой, как указано выше. Доступ к этому массиву «мьютексирован» с помощью атомарного int. (Я новичок в этом, так что, возможно, смешаю здесь несколько слов.)
Каждой рабочей группе неявно назначается часть массива входных элементов. Он использует локальный массив сегментов, содержащий новые индексы относительно начала (глобального) сегмента, о котором мы еще не знаем. После заполнения одного из этих «локальных буферов» рабочая группа должна записать локальные буферы в глобальный массив. Для этого он блокирует доступ к глобальному массиву начала сегмента, увеличивает эти значения на текущие размеры локального сегмента, разблокирует, а затем может записать результат в глобальный массив new_index (путем добавления соответствующего смещения). Этот процесс повторяется до тех пор, пока не будут обработаны все присвоенные элементы.
Возникают два вопроса:
Это хороший подход? Я знаю, что чтение и запись из / в глобальную память, скорее всего, является узким местом здесь, тем более что я пытаюсь получить синхронизированный доступ к (по крайней мере, только небольшой части ) глобальная память. Но, возможно, есть гораздо лучший подход для этого, возможно, используя декомпозицию ядра. Обратите внимание, что я стараюсь избегать обратного чтения данных из GPU в CPU во время ядер (чтобы избежать сброса очереди команд OpenCL, что тоже плохо, как меня учили).
В приведенном выше алгоритме как реализовать механизм блокировки? Будет ли работать что-то вроде следующего кода? В частности, я ожидаю проблем, когда оборудование выполняет рабочие элементы «по-настоящему параллельно» в группах SIMD, например, «перекосы» Nvidia. В моем текущем коде все элементы рабочей группы будут пытаться получить блокировку в режиме SIMD. Следует ли ограничить это только первым рабочим элементом? И использовать барьеры, чтобы синхронизировать их локально?
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable __kernel void putInBuckets(__global uint *mutex, __global uint *bucket_start, __global uint *new_index) { __local bucket_size[NUM_BUCKETS]; __local bucket[NUM_BUCKETS][LOCAL_MAX_BUCKET_SIZE]; // local "new_index" while (...) { // process a couple of elements locally until a local bucket is full ... // "lock" while(atomic_xchg(mutex, 1)) { } // "critical section" __local uint l_bucket_start[NUM_BUCKETS]; for (int b = 0; b < NUM_BUCKETS; ++b) { l_bucket_start[b] = bucket_start[b]; // where should we write? bucket_start[b] += bucket_size[b]; // update global offset } // "unlock" atomic_xchg(mutex, 0); // write to global memory by adding the offset for (...) new_index[...] = ... + l_bucket_start[b]; } }
