我正在阅读《大规模并行处理器编程》(第3版)这本书,书中介绍了Kogge-Stone并行扫描算法的实现。该算法旨在由单个块运行(这只是初步简化),下面是实现。
// X is the input array, Y is the output array, InputSize is the size of the input array
__global__ void Kogge_Stone_scan_kernel(float* X, float* Y, int InputSize) {
__shared__ float XY[SECTION_SIZE]; // SECTION_SIZE is basically blockDim.x
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < InputSize)
XY[threadIdx.x] = X[i];
for (unsigned int stride = 1; stride < blockDim.x; stride *= 2) {
__syncthreads();
if (threadIdx.x >= stride)
XY[threadIdx.x] += XY[threadIdx.x - stride]; // Race condition here?
}
Y[i] = XY[threadIdx.x];
}
不管算法的工作方式如何,我对XY[threadIdx.x] += XY[threadIdx.x - stride]
这行有点困惑。比如说stride = 1
,那么带有threadIdx.x = 6
的线程将执行XY[6] += XY[5]
操作。但是,同时具有threadIdx.x = 5
的线程将执行XY[5] += XY[4]
。问题是:是否可以保证线程6
将读取XY[5]
的原始值而不是XY[5] + XY[4]
?注意,这不限于其中锁步执行可以防止竞争条件的单个线程束。
谢谢
1条答案
按热度按时间neskvpey1#
是否保证线程6将读取XY[5]的原始值而不是XY[5] + XY[4]
不,CUDA不提供线程执行顺序的保证(锁步或其他方式),代码中也没有任何内容来解决这个问题。
顺便说一下,
cuda-memcheck
和compute-sanitizer
在识别共享内存竞争条件方面非常出色:正如你可能已经猜测到的,你可以通过在有问题的行中分解读和写操作来解决这个问题,在中间设置一个屏障: