为什么这段 cuda 的并行前缀算法不会数据竞争

rungo 2026-07-05 23:28 1

比如在 iter=1 这一轮的时候,同时会写入和读取 sum_buf[1]


naive_ker = SourceModule(""" 
__global__ void naive_prefix(double *vec, double *out)
{
__shared__ double sum_buf[1024];
int tid = threadIdx.x;
sum_buf[tid] = vec[tid];

int iter = 1;
for (int i=0; i < 10; i++)
{
__syncthreads();
if (tid >= iter )
{
sum_buf[tid]= sum_buf[tid] + sum_buf[tid - iter];
}
iter *= 2;
}
__syncthreads();
out[tid] = sum_buf[tid];
__syncthreads();
}
""")
naive_gpu = naive_ker.get_function("naive_prefix")
最新回复 (6)
  • r6cb 07-05 23:43
    1
    你这都是写入到 sum_buf[tid],每个线程 tid 不同,怎么会和其他线程冲突呢
  • rungo 楼主 07-05 23:52
    2
    @r6cb 比如 iter=1 时, tid=2 线程会读取 sum_buf[1] ,tid=1 线程会写入 sum_buf[1]
  • flyqie 07-06 00:35
    3
    好奇 AI 对这种问题的解释怎么样?

    感觉现在 AI 应该对这种还算"聪明"?
  • rungo 楼主 07-06 00:50
    4
    @flyqie 问了几个 ai 不能统一,这段代码我从书上抄过来的
  • r46mht 07-06 05:29
    5
    理论上可能会数据竞争,实际上数据竞争不太可能。如果你刻意的让一个 thread 等一等再算,就有可能看到 undfined 的结果:


    ```
    __global__ void naive_prefix(double *vec, double *out)
    {
    __shared__ double sum_buf[1024];
    int tid = threadIdx.x;
    sum_buf[tid] = vec[tid];

    int iter = 1;
    for (int i=0; i < 10; i++)
    {
    __syncthreads();

    // ARTIFICIAL DIVERGENCE:
    // Force even-numbered threads to stall for 500,000 clock cycles.
    // Odd-numbered threads will blow past this and execute the math below.
    if (tid % 2 == 0) {
    long long start = clock64();
    while(clock64() - start < 500000) {}
    }

    if (tid >= iter )
    {
    // The odd threads will update sum_buf before the even threads even reach this line.
    // When the even threads finally wake up, they will read corrupted, already-updated data.
    sum_buf[tid]= sum_buf[tid] + sum_buf[tid - iter];
    }

    iter *= 2;
    }
    __syncthreads();
    out[tid] = sum_buf[tid];
    }
    ```

    我猜 shared_memory 写入要比从 cache 读慢的多的多,并且 GPU thread 之间不像 CPU 一样有 cache coherence 的,所以直接读基本上不可能读到其他 thread 写入的结果。
  • lucays 07-06 14:19
    6
    跑了下 5.5 xhigh 的说法,这应该不算 AI 答复吧。。:
    这段代码确实有数据竞争;只是很多时候“看起来没错”,主要因为 GPU 的实际执行经常碰巧满足了你想要的顺序,尤其是同一个 warp 内。
    关键点:
    sum_buf 是 shared memory ,不是普通 cache 。它是 block 内所有 thread 共享的片上内存;一个 thread 写了,其他 thread 后续读是可以读到这个新值的。CUDA 不是因为“没有 cache coherence”所以读不到别的 thread 的写入。shared memory 的可见性和顺序要靠同步原语定义,尤其是 __syncthreads() 或 warp 内的 __syncwarp()。
    你原代码的问题是这一句:
    sum_buf[tid] = sum_buf[tid] + sum_buf[tid - iter];
    同一轮里,tid=1 会写 sum_buf[1],而 tid=2 会读 sum_buf[1]。没有第二个 barrier 把“所有线程先读旧值”和“所有线程再写新值”隔开,所以 CUDA 语义上这是 race 。
    为什么平时看不到的常见原因是:
    同一个 warp 内很多指令近似 lockstep 执行。
    编译后的指令通常是先 load ,再 add ,再 store 。
    如果一个 warp 的所有 lane 基本一起执行 load ,那么 tid=2 很可能在 tid=1 store 之前已经读完旧的 sum_buf[1]。
    小规模测试、固定输入、固定 GPU 架构上,调度顺序可能稳定,所以错误不暴露。
    但这只是实现细节上的“碰巧”,不是 CUDA 保证。

    问了下可能正确的代码:
    ```
    for (int iter = 1; iter < 1024; iter *= 2) {
    __syncthreads();

    double x = 0.0;
    if (tid >= iter) {
    x = sum_buf[tid - iter];
    }

    __syncthreads();

    if (tid >= iter) {
    sum_buf[tid] += x;
    }
    }
    ```
* 帖子来源V2EX
返回