Проблемы с общей памятью (?) В CUDA (JCUDA) / неопределенное поведение - PullRequest
Новая
       26

Проблемы с общей памятью (?) В CUDA (JCUDA) / неопределенное поведение

1 голос
/ 16 апреля 2011

Я работаю над своим игровым проектом (защита башни) и пытаюсь вычислить расстояние между всеми критериями и башней с помощью JCuda, используя общую память.Для каждой башни я запускаю 1 блок с N threds, где N равно количеству тварей на карте.Я вычисляю расстояние между всеми критериями и этой башней для данного блока, и я храню наименьшее найденное расстояние в общей памяти блока.Мой текущий код выглядит так: 

extern "C"

__global__ void calcDistance(int** globalInputData, int size, int
critters, int** globalQueryData, int* globalOutputData) {

 //shared memory
 __shared__ float minimum[2];

 int x = threadIdx.x  + blockIdx.x * blockDim.x;
 int y = blockIdx.y;

 if (x < critters) {

   int distance = 0;
   //Calculate the distance between tower and criter
   for (int i = 0; i < size; i++) {
     int d = globalInputData[x][i] - globalQueryData[y][i];
     distance += d * d;
   }

   if (x == 0) {        
     minimum[0] = distance;
     minimum[1] = x;
   }

   __syncthreads();



   if (distance < minimum[0]) {
     minimum[0] = distance;
     minimum[1] = x;
   }

   __syncthreads();
   globalOutputData[y * 2]     = minimum[0];
   globalOutputData[y] = minimum[1];

 }


}

Проблема заключается в том, что я повторно запускаю код, используя один и тот же ввод несколько раз (я освобождаю всю память на хосте и устройстве после каждого запускаКаждый раз, когда код выполняется для числа блоков (башни)> 27, я получаю разные выходные данные. Я совершенно уверен, что это как-то связано с разделяемой памятью и с тем, как я им занимаюсь, - переписывая код для использования globalпамять дает тот же результат всякий раз, когда код выполняется.Есть идеи?

1 Ответ

1 голос
/ 16 апреля 2011

В этом ядре есть проблема с гонкой памяти (поэтому правильность чтения после записи): 

   if (distance < minimum[0]) {
     minimum[0] = distance;
     minimum[1] = x;
   }

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

Более мягкая версия той же проблемы относится к обратной записи в глобальную память в конце ядра: 

   __syncthreads();
   globalOutputData[y * 2]     = minimum[0];
   globalOutputData[y] = minimum[1];

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

Если ваше намерение состоит в том, чтобы каждый поток вычислял расстояние, а затем для минимального значения расстояния по блоку для записи в глобальную память, вам придется либо использовать атомарные операции с памятью (для разделяемой памяти это поддерживается только на устройствах Compute 1.2 / 1.3 и 2.x), либо записывать явное сокращение общей памяти. После этого только один поток должен выполнить обратную запись в глобальную память.

Наконец, у вас также есть потенциальная проблема с корректностью синхронизации, которая может привести к зависанию ядра. __syncthreads() (который отображается на инструкцию панели PTX) требует, чтобы каждый поток в блоке прибыл и выполнил инструкцию до продолжения ядра. Имея такой поток управления: 

 if (x < critters) {
 ....
   __syncthreads();
 ....
 }

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

Итак, в заключение вернемся к чертежной доске по крайней мере по трем вопросам в текущем коде.

...