对有多个线程写入相同索引的CUDA内核进行编码？

我正在编写一些代码来激活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是连接i的源神经元索引，以及连接i的目标神经元索引。 weightArray包含重量信息（所以索引i的weightArray对应于连接i）。

正如你所看到的，SumArray是我存储的总和。因此，内核将sumArray（连接目标神经元索引i）增加连接权重的绝对值i。直观地说，对于到神经元的所有传入连接，求和所有权重。这就是我正在试图用这个内核做的所有事情。最终，我会用这个总和来标准化权重。

问题是这是错误的。我已经连续做了这个，答案是不同的。答案不同，通常约12-15倍（所以正确的答案是700.0，而我得到的是50年代的东西）。

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

编辑：为了参考，__threadfence() usaged在CUDA编程指南V3.1附录B.5存储器防护功能

来源

2010-09-14 Paul

你需要做一个减少。

萨姆的元素分配给每个线程，并将结果存放在一个数组中，高速缓存[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

此示例代码是她E：

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

它也相当不错，在“CUDA通过实施例”说明（这是其中代码片段来自）。

这种方法有一个很大的警告。这些添加不会以与串行代码相同的顺序发生。浮动的添加不可交换，因此舍入错误可能会导致稍微不同的结果。

来源

2010-10-01 06:19:13

在上面的示例代码中，'cache [cacheIndex] + 1'应该是'cache [cacheIndex + 1]'。 – 2016-11-25 06:14:22

+=描述的内容不是atomical =>不是线程安全。使用atomicAdd。

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

来源

2010-09-14 16:52:32 Andrey

我不知道我明白。 atomicAdd是用于整数，我使用浮动。另外，当你说“单独计算和合并”时，我的场景中引用的单个计算是什么？总和？我不知道如何避免写入同一个单元格。 – Paul 2010-09-14 17:04:05

@Paul Open B.11.1.1 NVIDIA CUDA C编程指南3.1版 5/28/2010。有atomicAdd的浮动版本。好的，你的情况你没有单独的计算。你写的代码效率不高。请阅读有关如何有效总结的更多信息：http://http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html – Andrey 2010-09-14 17:44:58

AtomicAdd支持浮动，但仅限于以后的CUDA版本。在CUDA 2.0之前，只支持整数AtomicAdd。 – 2010-10-01 06:05:54

对有多个线程写入相同索引的CUDA内核进行编码？

回答

相关问题