CUDA множественный доступ к памяти

Question

Пожалуйста, объясните мне, как работает доступ к памяти в следующем ядре:

__global__ void kernel(float4 *a)
{
     int tid = blockIdx.x * blockDim.x + threadIdx.x;

     float4 reg1, reg2;
     reg1 = a[tid]; //each thread reads a unique memory location

     for(int i = 0; i < totalThreadsNumber; i++)
     {  
          reg2 = a[i]; //all running threads start reading 
                       //the same global memory location
          //some computations
     }

     for(int i = 0; i < totalThreadsNumber; i++)
     {
          a[i] = reg1; // all running threads start writing 
                       //to the same global memory location
                       //race condition
     }
}

Как это работает в первом цикле?Есть ли сериализация?Я предполагаю, что второй цикл вызывает сериализацию потоков (только внутри деформации?) И результат не определен.

Answer 1

В соответствии с моим объяснением Ферми (sm_2x), на старом оборудовании доступ к памяти осуществляется за полвертыва.

В первом цикле (чтение) вся деформация считывает с того же адреса в локальную переменную. Это приводит к «трансляции». Поскольку Fermi имеет кэш L1, либо будет загружена одна строка кэша, либо данные будут извлечены непосредственно из кэша (для последующих итераций). Другими словами, нет сериализации.

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