r46mht

理论上可能会数据竞争，实际上数据竞争不太可能。如果你刻意的让一个 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 写入的结果。

其实还是更喜欢露营睡大觉 lol

刑不可知 则威不可测

Claude 打击违规用户跟我精通中文的尼日利亚天才程序员 Musa 有什么关系

放轻松，出事也不一定是在天上

[突发！美国纽约市拉瓜迪亚机场发生飞机与消防车相撞事故，造成多人伤亡，目前报道至少有 2 人死亡，约 70 人受伤] https://www.bilibili.com/video/BV1XoQdBGERR/

你打反欺诈部门的人工电话了吗？没打打打试试呗

@mayday1997

> 中国滑雪规范明确提出了转弯让直行

你们律师上庭的核心论点不会是这个吧...