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

Posted

技术标签:

【中文标题】编写一个有许多线程写入同一索引的 CUDA 内核?【英文标题】:Coding a CUDA Kernel that has many threads writing to the same index? 【发布时间】:2011-04-12 06:44:06 【问题描述】:

我正在编写一些代码来激活 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 Programming Guide v3.1 附录 B.5 Memory Fence Functions 中描述了使用的__threadfence()

【问题讨论】:

【参考方案1】:

你需要做一个减少。

对分配给每个线程的元素求和并将结果放入一个数组中,cache[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”中也有很好的解释(这是代码片段的来源)。

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

【讨论】:

在上面的示例代码中,cache[cacheIndex] + 1 应该是cache[cacheIndex+1]【参考方案2】:

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

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

【讨论】:

我不确定我是否理解。 atomicAdd 用于整数,我使用浮点数。此外,当您说“单独计算和合并”时,我的场景中引用的单独计算是什么?总和?我不确定如何避免写入同一个单元格。 @Paul Open B.11.1.1 of NVIDIA CUDA C Programming Guide Version 3.1 5/28/2010。 atomicAdd 有float 版本。好的,在您的情况下,您没有单独的计算。你写的代码效率不高。在此处阅读有关如何有效求和的更多信息:http.developer.nvidia.com/GPUGems3/gpugems3_ch39.html AtomicAdd 支持浮点数,但仅在更高的 CUDA 版本上。在 CUDA 2.0 之前,仅支持整数 AtomicAdd。

以上是关于编写一个有许多线程写入同一索引的 CUDA 内核?的主要内容,如果未能解决你的问题,请参考以下文章

多个线程在同一个全局内存位置写入相同的值

CUDA 中大小为 4 的非法写入

通过表面写入 CUDA 中的浮点 OpenGL 纹理

CUDA 中的分歧 - 从内核中的线程退出

从同一位置读取时 CUDA 中的银行冲突

同一 GPU 上的 OpenCL 和 CUDA 内核