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

保留我对 Fermi (sm_2x) 的解释，在较旧的硬件上，内存访问是按半扭曲进行的。

在第一个循环（读取）中，整个扭曲从同一地址读取到局部变量中。这导致“广播”。由于 Fermi 有一个 L1 缓存，要么会加载一个缓存行，要么直接从缓存中获取数据（用于后续迭代）。换句话说，没有序列化。

在第二个循环（写入）中，哪个线程获胜是未定义的 - 就像任何多线程编程模型一样，如果多个线程写入同一位置，则程序员负责理解竞争条件。您无法控制块中的哪个扭曲将最后执行，也无法控制最后一个扭曲中的哪个线程将完成写入，因此您无法预测最终值是什么。