CUDA: atomicAdd занимает слишком много времени, сериализуя потоки

У меня есть ядро, которое делает некоторые сравнения и решает, сталкиваются ли два объекта или нет. Я хочу сохранить идентификаторы сталкивающихся объектов в выходной буфер. Я не хочу иметь пробел в буфере вывода. Я хочу записать каждое столкновение с уникальным индексом в буфере вывода.

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

__global__ void mykernel(..., unsigned int *gColCnt) {
    ...

    __shared__ unsigned int sColCnt;
    __shared__ unsigned int sIndex;

    if (threadIdx.x == 0) {
        sColCnt = 0;
    }

    __syncthreads();

    unsigned int index = 0;
    if (colliding)
        index = atomicAdd(&sColCnt, 1); //!!Time Consuming!!

    __syncthreads();

    if (threadIdx.x == 0)
        sIndex = atomicAdd(gColCnt, sColCnt);

    __syncthreads();

    if (sColCnt + sIndex > outputSize) { //output buffer is not enough
        //printf("Exceeds outputsize: %d + %d > %d\n", sColCnt, sIndex, outputSize);
        return;
    }

    if (colliding) {
        output[sIndex + index] = make_uint2(startId, toId);
    }
}

Моя проблема в том, что, когда многие потоки пытаются увеличить атомарную переменную, они сериализуются. Прежде чем написать что-то вроде prefix-sum, я хотел спросить, есть ли способ сделать это эффективно.

Истекшее время моего ядра увеличивается с 13 мс до 44 мс из-за этой строки.

Я нашел пример кода с префиксной суммой, но ссылки на него не работают из-за неработающей доски обсуждений NVIDIA. /questions/7896873/kodirovanie-yadra-cuda-s-mnozhestvom-potokov-zapisyivayuschih-v-odin-i-tot-zhe-indeks/7896880#7896880

Изменить: я добавил конец моего кода тоже выше. На самом деле у меня есть иерархия. Чтобы увидеть влияние каждой строки кода, я настраиваю сцены, где каждый объект сталкивается друг с другом, экстремальный случай и другой экстремальный случай, когда объекты практически не сталкиваются.

В конце я добавляю общую атомарную переменную в глобальную переменную (gColCnt), чтобы сообщать извне о количестве столкновений и находить правильные значения индекса. Я думаю, что я должен использовать AtomicAdd здесь в любом случае.

Источник

2 ответа

Решение

Рассмотрите возможность использования алгоритма сжатия параллельного потока, например thrust::copy_if,

Источник

Статья в блоге nvidia по теме: http://devblogs.nvidia.com/parallelforall/gpu-pro-tip-fast-histograms-using-shared-atomics-maxwell/

Источник
Другие вопросы по тегам