0

有人可以帮我理解以下CUDA-C的部分并行求和算法实现吗？我很难理解共享partialSum数组的初始填充 [第 3 到 8 行]。我已经追踪了几个小时，但我不明白为什么应该从下面的代码开始而2*blockIdx.x*blockDim.x;不是blockIdx.x*blockDim.x;？

主机代码：

numOutputElements = numInputElements / (BLOCK_SIZE<<1);
 if (numInputElements % (BLOCK_SIZE<<1)) {
     numOutputElements++;
 }
#define BLOCK_SIZE 512
dim3 dimGrid(numOutputElements, 1, 1);
dim3 dimBlock(BLOCK_SIZE, 1, 1);
total<<<dimGrid, dimBlock>>>(deviceInput, deviceOutput, numInputElements);

内核代码：

1    __global__ void total(float * input, float * output, int len) {
2    
3    __shared__ float partialSum[2*BLOCK_SIZE];
4      
5      unsigned int t = threadIdx.x;
6      unsigned int start = 2*blockIdx.x*blockDim.x;
7      partialSum[t] = input[start + t];
8      partialSum[blockDim.x + t] = input[start + blockDim.x + t];
9     
10    for (unsigned int stride = blockDim.x; stride >=1; stride >>=1)
11      {
12       __syncthreads();
13        
14       if (t < stride)
15         partialSum[t] += partialSum[t + stride];
16      }
17      output[blockIdx.x] = partialSum[0];   
18  }

假设我有 10 个元素要求和，我选择将块大小设为 4，每个块有 4 个线程，所以将有 3 个块在使用，对吧？[让我们暂时忘记经线大小和其他事情]

当 blockIdx.x 为 2 （具有 2 个元素的最后一个块）时，开始变为 (2*2*4=)16 并且大于 10 并且超过input长度（因此partialSum[t]和partialSum[blockDim.x + t]都将保持不变并且block2's共享内存将保持为空。）如果是这样，那么我数组的最后 2 个元素将丢失！

这让我觉得我得到了 blockIdx.x, blockDim.x 错误的方式。有人可以纠正我吗？请！

4

2 回答 2

3

你只会启动一半的块，每个块做两倍的工作。这样做的好处是存储部分和所需的暂存空间减少了一半（因为您只启动了一半的块）。

进行减少（在这种情况下为总和）的通常方法是做这样的事情。

1    __global__ void total(float * input, float * output, int len) {
2    
3    __shared__ float partialSum[BLOCK_SIZE];
4      
5     unsigned int t = threadIdx.x;
6     unsigned int start = blockIdx.x*blockDim.x;
7     partialSum[t] = 0;
8     for (int T = start; T < len; T += blockDim.x * gridDim.x) 
9        partialSum[t] += input[T];
10    for (unsigned int stride = blockDim.x/2; stride >=1; stride >>=1)
11      {
12       __syncthreads();
13        
14       if (t < stride)
15         partialSum[t] += partialSum[t + stride];
16      }
17      output[blockIdx.x] = partialSum[0];   
18  }

所以如果你有len = 1024, 和BLOCK_SIZE = 256，你可以启动任何 <= 4 块。

让我们看看当你启动不同数量的块时，第 8 行和第 9 行中包含的 for 循环会发生什么。还要记住输出需要有元素的数量==块的数量。

  • Blocks == 4意味着，blockDim.x * gridDim.x= 256 x 4 = 1024，所以它只会迭代一次。对输出的非合并写入次数 = 4。
  • Blocks == 2意思是，blockDim.x * gridDim.x= 256 x 2 = 512，所以它会迭代两次。对输出的非合并写入次数 = 2。
  • Blocks == 1意思是，blockDim.x * gridDim.x= 256 x 1 = 256，所以它会迭代 4 次。对输出的非合并写入次数 = 1。

因此，启动更少的块有利于减少内存占用并减少全局写入。然而，它降低了并行性。

理想情况下，您需要启发式地找到最适合您的算法的组合。或者你可以使用现有的库来为你做这件事。

有问题的内核选择启动一半的块以获得一些性能改进。但是可能不需要使用两倍的共享内存。

于 2013-01-12T07:38:25.840 回答
1

您可能在计算积木时遇到问题。

假设您有 10 个元素要求和，并且您选择将块大小设为 4，并且每个块有 4 个线程，那么将只有两个块在使用中。

根据您的内核代码，由于每个线程负责全局设备内存中的两个元素。

每个线程读取的输入元素如下所示。我没有在您的代码中看到任何范围检查。所以我假设你的 10 个元素有足够的零填充。

blockIdx.x           : 0 0 0 0  1 1 1 1  2 2 2 2  3 3 3 3
threadIdx.x          : 0 1 2 3  0 1 2 3  0 1 2 3  0 1 2 3
linear thread id     : 0 1 2 3  4 5 6 7  8 9 a b  c d e f

Idx of the element     0 1 2 3  8 9
read by the thread   : 4 5 6 7

所以output[0]存储 elmemnt 0~7output[1]的总和，存储元素 8~9 的总和。我不认为有什么损失。

请参阅优化 CUDA 中的并行减少中的内核 4，以了解有关为什么存在2*. @Pavan在他的回答中给出的较慢的内核 3 和内核是类似的实现，其中每个线程只负责一个元素。

于 2013-01-12T12:26:51.803 回答