Вопрос
Предположим, несколько рабочих элементов хотят добавить в глобальный стек:
void kernel(__global int* stack) {
... do stuff ...
push(stack, value);
... do stuff ...
return y;
}
Желательно, чтобы после запуска ядра stack содержал каждое valueподтолкнул к этому.Заказ не имеет значения.Как правильно сделать это в OpenCL 1.2?
Что я пробовал
Очевидная идея - использовать atomic_inc, чтобы получить длину и просто написать в нее:
void push(__global int* stack, int val) {
int idx = atomic_inc(stack) + 1; // first element is the stack length
stack[idx] = val;
}
Но я предполагаю, что все рабочие элементы, вызывающие atomic_inc отдельно в одной и той же позиции памяти, разрушают параллелизм.Отдельной идеей было бы просто записать во временный массив больше, чем количество рабочих элементов:
void push(__global int* stack, int val) {
stack[get_global_id(0)] = val;
}
Это оставило бы нас с редким массивом значений:
[0, 0, 0, 7, 0, 0, 0, 2, 0, 0, 3, 0, 0, 0, 9, 0, 0, ...]
Который затем может быть сжат с помощью «потоковое сжатие» .Поэтому мне интересно, какая из этих идей наиболее эффективна, и, возможно, есть третий вариант, о котором я не знаю.