编写具有多个线程写入同一索引的CUDA内核？
我正在写一些在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是connection i的源神经元索引，也是connection i的目标神经元索引。weightArray包含权重信息(因此weightArray的索引i对应于连接i)。
如您所见，SumArray是我存储总和的位置。因此内核将sumArray (在连接i的目标神经元索引处)递增连接i的权重的绝对值。直观地说，对于到神经元的所有传入连接，将所有权重相加。这就是我想要用这个内核做的所有事情。最后，我将使用这个和来归一化权重。
问题是它是错误的。我已经连续地这样做了，答案是不同的。答案是不同的，通常是12-15倍(所以正确的答案是700.0，我得到的是50s的范围)。
您可以看到，我添加了__threadfence() (和__threadfence_block()，以确保写入不是由每个线程在同一时间完成)。我不确定这是否是我代码的问题所在。我已经确保权重数组与我测试的串行版本相同，并且源/目标信息也是相同的。我做错了什么？
编辑:作为参考，CUDA编程指南v3.1附录B.5内存围栏功能中描述了__threadfence()使用的功能
回答 2
Stack Overflow用户
发布于 2010-10-01 06:19:13
你需要做一个消减。
对分配给每个线程的元素求和，并将结果放入一个数组中，然后使用cachethreadsPerBlock和__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“中也有很好的解释(这就是代码片段的来源)。
这种方法有一个很大的警告。添加将不会以与序列码相同的顺序发生。浮点数的相加是不可交换的，因此舍入误差可能会导致略有不同的结果。
Stack Overflow用户
发布于 2010-09-14 16:52:33
+=不是原子的=>，不是线程安全的。使用atomicAdd。
此外，您还应避免写入相同的存储单元。问题是这些调用将被序列化，线程将排队等待。如果你无法避免这个操作，试着把你的算法分成两个阶段:单独计算和合并。并行合并可以非常高效地实现。
https://stackoverflow.com/questions/3711027
复制相似问题
腾讯云开发者
