CUDA。如何展开前32个线程，以便它们并行执行？

我知道“每个warp包含连续的，增加线程ID的线程，第一个warp包含线程0”，所以前32个线程应该在第一个warp中。另外我知道一个warp中的所有线程都可以在任何可用的Streaming Multiprocessor上同时执行。CUDA。如何展开前32个线程，以便它们并行执行？

据我所知，因为如果只有一个warp正在执行，没有必要在线程同步。但是，如果我删除倒数第二个if块中的任何__syncthreads()块，则下面的代码会产生错误的答案。我试图找到原因，但没有结果。我真的希望得到你的帮助，所以你可以告诉我这段代码有什么问题？为什么我不能只留下最后的__syncthreads()并得到正确答案？

#define BLOCK_SIZE 128 

__global__ void reduce (int * inData, int * outData) 
{ 
__shared__ int data [BLOCK_SIZE]; 
int tid = threadIdx.x; 
int i = blockIdx.x * blockDim.x + threadIdx.x; 

data [tid] = inData [i] + inData [i + blockDim.x/2 ]; 
__syncthreads(); 

for (int s = blockDim.x/4; s > 32; s >>= 1) 
{ 
    if (tid < s) 
    data [tid] += data [tid + s]; 
    __syncthreads(); 
} 

if (tid < 32) 
{ 
    data [tid] += data [tid + 32]; 
    __syncthreads(); 
    data [tid] += data [tid + 16]; 
    __syncthreads(); 
    data [tid] += data [tid + 8]; 
    __syncthreads(); 
    data [tid] += data [tid + 4]; 
    __syncthreads(); 
    data [tid] += data [tid + 2]; 
    __syncthreads(); 
    data [tid] += data [tid + 1]; 
    __syncthreads(); 
} 
if (tid == 0) 
    outData [blockIdx.x] = data [0]; 
} 

void main() 
{ 
... 
reduce<<<dim3(128), dim3(128)>>>(dev_data, dev_res); 
... 
}

P.S.我使用GT560Ti

来源

2012-12-21 Виталий Хабаров

你应该申报的共享内存变量挥发性：

__shared__ volatile int data [BLOCK_SIZE];

你所看到的问题是Fermi架构和编译器优化的神器。费米架构缺乏直接操作共享内存的指令（它们出现在G80/90/GT200系列中）。因此，所有内容都被加载到注册，操作并存储回共享内存。但是编译器可以自由地推断出，如果一系列操作在寄存器中进行，没有中间加载和存储共享内存，代码可以变得更快。这是非常好的，除了，当你在相同的warp操作共享内存中依赖隐式同步线程时，就像这种简化代码一样。

通过声明共享内存缓冲区为volatile，迫使编译器在每个缩减阶段之后强制执行共享内存写入，并且还原了warp内的线程之间的隐式数据同步。

此问题在费米的编程笔记中讨论，该笔记将随CUDA工具包一起发运（或可能附带）。

来源

2012-12-21 17:07:00 talonmies

CUDA。如何展开前32个线程，以便它们并行执行？

回答

相关问题