Cuda - выборочный магазин памяти

Question

В моем ядре, если условие выполнено, я обновляю элемент буфера вывода

if (condition(input[i])) //?
    output[i] = 1;

в противном случае вывод может остаться прежним, имея значение 0.

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

У меня вопрос: лучше написать все элементы, добиться объединения или сделать выборочную запись?

output[i] = condition(input[i]); //? 

Не могли бы вы обсудить ваши заявления?

Answer 1

Объединение достигается, даже если некоторые потоки в основе не участвуют в загрузке или хранении, если все участвующие потоки удовлетворяют требованиям объединения. Поэтому условные записи не должны влиять на пропускную способность памяти.

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

Answer 2

На моем установочном ядре, которое делает условный набор (вариант 1), работает 1,77 доллара США, а вариант 2 1,399 доллара США. Это мой код (setConditional - более быстрый):

__global__ void conditionalSet(unsigned int* array) {
    if ((threadIdx.x & 3) == 0) {
        array[threadIdx.x] = 1;
    }
}

__global__ void setConditional(unsigned int* array) {
    array[threadIdx.x] = (threadIdx.x & 3) == 0 ? 1 : 0;
}