在共享内存中移动数据时的 CUDA 线程执行

Question

我有以下功能：

代码示例 1：

__global__ void func(const int *input, int N){

  extern __shared__int buffer[];
  int temp = 0;

  for(int i = blockIdx.x*blockDim.x + threadIdx.x; i < N; i += blockDim.x*gridDim.x; ){
     temp += input[i];
  }

  buffer[threadIdx.x] = temp;
  __syncthreads();

} 

它是并行归约函数的一部分。据我了解，它从全局内存复制到共享内存。

我试图通过一个简单的例子来理解它。例如我有一个大小为 20 的一维数组 元素（N = 20）。我想象执行如下。纠正我如果我错了。对于 5 个块，每个块 4 个线程。

第一个块的所有线程执行：

blockIdx.x=0
threadIdx.x=0
for(i=0; i<18; i+= 4*5){ temp= in[0] /i wrote the sums intuitively/}        
buffer[threadIdx.x] = temp

blockIdx.x=0
threadIdx.x=1
for(i=1; i<18; i+= 4*5){ temp= in[1] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

blockIdx.x=0
threadIdx.x=2
for(i=2; i<18; i+= 4*5){ temp= in[2] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

blockIdx.x=0
threadIdx.x=3
for(i=3; i<18; i+= 4*5){ temp= in[3] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

执行第二块的所有线程：

blockIdx.x=1
threadIdx.x=0
for(i=1*4; i<18; i+= 4*5){ temp= in[4] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

blockIdx.x=1
threadIdx.x=1
for(i=1*4+1; i<18; i+= 4*5){ temp = in[5] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

blockIdx.x=1
threadIdx.x=2
for(i=1*4+2; i<18; i+= 4*5){ temp = in[6] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

blockIdx.x=1
threadIdx.x=3
for(i=1*4+3; i<18; i+= 4*5){ temp = in[7] /i wrote the sums intuitively/}
buffer[threadIdx.x] = temp

等

为什么我们有一个for循环而不是写：

代码示例 2：

unsigned int i = blockIdx.x*blockDim.x + threadIdx.x;
buffer[threadIdx.x] = input[i]; 

谁能给出一个直观的例子或解释？

Answer 1

在主函数中，索引大于blockDim.x*(gridDim.x-1)+(blockDim.x-1)的元素将被考虑在计算中，而在您提供的方法中则不会发生。

假设您有N=1024，并且您使用一个网格调用您的函数，该网格有 8 个块，每个块有 32 个线程。在您的主函数中，线程i 将在元素i、i+8*32、i+2*(8*32)、i+3*(8*32) 处收集并累加属于*input 的数据。另一方面，您的代码仅在元素 i 处收集数据。换句话说，它仅将*input 的第一个元素32*8 相加，而忽略1024-32*8 其余元素。

更详细：

代码示例 1 的工作方式如下：

blockIdx.x=0
threadIdx.x=0
for ( i = 0; i < 1024; i += 32*8 )
    temp += input[i]; // temp= input[0]+input[256]+input[512]+input[768]
buffer[0] = temp; //=input[0]+input[256]+input[512]+input[768]

blockIdx.x=0
threadIdx.x=1
for ( i = 1; i < 1024; i += 32*8 )
    temp += input[i]; // temp= input[1]+input[257]+input[513]+input[769]
buffer[1] = temp; //=input[1]+input[257]+input[513]+input[769]

blockIdx.x=0
threadIdx.x=2
for ( i = 2; i < 1024; i += 32*8 )
    temp += input[i]; // temp= input[2]+input[258]+input[514]+input[770]
buffer[2] = temp; //=input[2]+input[258]+input[514]+input[770]

...

//last thread
blockIdx.x=7
threadIdx.x=31
for ( i = 7*32+31; i < 1024; i += 32*8 )
    temp += input[i]; // temp= input[255]+input[511]+input[767]+input[1023]
buffer[255] = temp; //=input[255]+input[511]+input[767]+input[1023]

代码示例 2 的工作方式如下：

blockIdx.x=0
threadIdx.x=0
i = 0*32+0; //=0
buffer[0] = input[0];

blockIdx.x=0
threadIdx.x=1
i = 0*32+1; //=1
buffer[1] = input[1];

....

//last thread
blockIdx.x=7
threadIdx.x=31
i = 7*32+31; //=255
buffer[255] = input[255];

如您所见，第一个代码示例遍历 input 数组的所有元素，但第二个代码示例没有。