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

Question

Я пишу код для активации нейронных сетей в 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 Функции забора памяти

Answer 1

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

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

Answer 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» (отсюда и фрагмент кода).

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