编写一个有多个线程写入同一索引的 CUDA 内核？

Question

我正在编写一些用于激活 CUDA 上的神经网络的代码，但遇到了问题。我没有得到进入给定神经元的权重的正确总和。

这是内核代码，我将尝试使用变量对其进行更清晰的解释。

__global__ void kernelSumWeights(float* sumArray, float* weightArray, int2* sourceTargetArray, int cLength)
{
int nx = threadIdx.x + TILE_WIDTH*threadIdx.y;
int index_in = (blockIdx.x + gridDim.x*blockIdx.y)*TILE_WIDTH*TILE_WIDTH + nx;
 if(index_in < cLength)
 {

  sumArray[sourceTargetArray[index_in].y] += fabs(weightArray[index_in]);
  //__threadfence();
  __threadfence_block();

 }

}

首先，网络中的连接数是cLength。对于每个连接，都有一个源神经元和一个目标神经元，以及该连接的权重。 SourceTargetArray 包含该信息。因此sourceTargetArray的索引i是连接i的源神经元索引，以及连接i的目标神经元索引。 weightArray 包含权重信息（因此 weightArray 的索引 i 对应于连接 i）。

如您所见，SumArray 是我存储总和的位置。因此，内核将 sumArray （在连接 i 的目标神经元索引处）增加连接 i 权重的绝对值。直观地说，对于神经元的所有传入连接，将所有权重相加。这确实是我想用这个内核做的全部事情。最终，我将使用这个总和来标准化权重。

问题是它是错误的。我连续这样做了，答案是不同的。答案有所不同，通常约为 12-15 倍（因此正确答案为 700.0，而我得到的结果在 50 倍范围内）。

您可以看到我添加了 __threadfence() （和 __threadfence_block() 试图确保写入不会由每个线程同时完成） 。我不确定这是否是我的代码的问题。我已确保权重数组与我测试的串行版本相同，并且源/目标信息也相同。我做错了什么？

编辑：作为参考，CUDA 编程指南 v3.1 附录 B.5 内存栅栏函数中描述了所使用的 __threadfence()

Answer 1

+= 不是原子的 =>不是线程安全的。使用atomicAdd。

此外，您还应该避免写入同一存储单元。问题是这些调用将被序列化，线程将排队并相互等待。如果您无法避免此操作，请尝试将算法分为两个阶段：单独计算和合并。并行合并可以非常有效地实现。

Answer 2

你需要做减量。

将分配给每个线程的元素相加，并将结果放入数组中，先缓存 [threadsPerBlock]，然后 __Syncthreads

现在，通过添加连续的相邻小计来减少结果小计：

int cacheIndex = threadIdx.x;
int i = blockDim.x / 2;
while (i != 0)
{
    if (cacheIndex < i)
        cache[cacheIndex] += cache[cacheIndex] + 1;
        __syncthreads;
        i /= 2;
    }
}

下面的内容详细解释了这一点：

http://developer.download.nvidia.com/compute/cuda/1_1 /Website/projects/reduction/doc/reduction.pdf

示例代码位于：

http://www.nvidia.com/object/cuda_sample_data-parallel.html

在“CUDA BY Example”（这是代码片段的来源）中也有很好的解释。

这种方法有一个很大的警告。添加的顺序不会与序列代码中的顺序相同。浮点数的相加不可交换，因此舍入误差可能会导致结果略有不同。