Кодирование ядра CUDA с множеством потоков, записывающих в один и тот же индекс?

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

Итак, вот код ядра, и я попытаюсь объяснить его с помощью переменных немного яснее.

__global__ void kernelSumWeights(float* sumArray, float* weightArray, int2* sourceTargetArray, int cLength)
{
int nx = threadIdx.x + TILE_WIDTH*threadIdx.y;
int index_in = (blockIdx.x + gridDim.x*blockIdx.y)*TILE_WIDTH*TILE_WIDTH + nx;
 if(index_in < cLength)
 {

  sumArray[sourceTargetArray[index_in].y] += fabs(weightArray[index_in]);
  //__threadfence();
  __threadfence_block();

 }

}

Во-первых, количество подключений в сети cLength, Для каждого соединения есть исходный нейрон и целевой нейрон, а также вес для этого соединения. SourceTargetArray содержит эту информацию. Итак индекс i из sourceTargetArray является исходным нейронным индексом соединения iи целевой нейронный индекс соединения i, weightArray содержит информацию о весе (так индекс i из weightArray соответствует соединению i).

Как вы видете, SumArray где я храню суммы. Так что ядро ​​увеличивает sumArray (у целевого нейронного индекса связи i) по абсолютной величине веса соединения i, Интуитивно, для всех входящих соединений с нейроном, суммируйте все веса. Это действительно все, что я пытаюсь сделать с этим ядром. В конце концов, я нормализую вес, используя эту сумму.

Проблема в том, что это неправильно. Я сделал это поочередно, и ответ другой. Ответы различаются, как правило, примерно в 12-15 раз (поэтому правильный ответ будет 700.0, и я получаю что-то в диапазоне 50-х).

Вы можете видеть, что я добавил __threadfence() (а также __threadfence_block() в попытке удостовериться, что записи не были выполнены одновременно каждым потоком). Я не уверен, что это проблема с моим кодом. Я гарантировал, что массив весов идентичен серийной версии, которую я тестировал, и что информация об источнике / цели также идентична. Что я делаю неправильно?

РЕДАКТИРОВАТЬ: для справки, __threadfence() Использование описано в Руководстве по программированию CUDA v3.1, Приложение B.5 Функции забора памяти

Источник

2 ответа

Решение

Вам нужно сделать сокращение.

Суммируйте элементы, назначенные каждому потоку, и поместите результат в массив, кэшируйте [threadsPerBlock], затем __Syncthreads

Теперь уменьшите итоговые промежуточные итоги, добавив последовательные соседние промежуточные итоги:

int cacheIndex = threadIdx.x;
int i = blockDim.x / 2;
while (i != 0)
{
    if (cacheIndex < i)
        cache[cacheIndex] += cache[cacheIndex] + 1;
        __syncthreads;
        i /= 2;
    }
}

Следующая колода объясняет это в некоторых деталях:

http://developer.download.nvidia.com/compute/cuda/1_1/Website/projects/reduction/doc/reduction.pdf

Пример кода для этого здесь:

http://www.nvidia.com/object/cuda_sample_data-parallel.html

Это также очень хорошо объяснено в "Примере CUDA BY" (отсюда и фрагмент кода).

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

Источник

+= не атомарный => не потокобезопасен. Используйте atomicAdd.

Также вам следует избегать записи в одну и ту же ячейку памяти. Проблема в том, что эти вызовы будут сериализованы, потоки будут стоять в очереди и ждать друг друга. Если вы не можете избежать этой операции, попробуйте разбить ваш алгоритм на две фазы: индивидуальные вычисления и слияние. Параллельное объединение может быть реализовано очень эффективно.

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