lucays's recent timeline updates
lucays

lucays

V2EX member #257375, joined on 2017-09-30 20:28:38 +08:00
Today's activity rank 974
Per lucays's settings, the topics list is hidden
Deals info, including closed deals, is not hidden
lucays's recent replies
我用普通鼠标才会手痛,换垂直鼠标就好了
键盘没感觉过啊

以前刚用 mac 那几天没有鼠标的时候 8 小时触摸板倒是手指痛过
10h 27m ago
Replied to a topic by 0x114514 生活 这个世界太割裂了
其实我们快乐教育有一个最大的问题
我们体力劳动没老外那么赚钱。。。
其实换成男的,零氪玩游戏也会被骂才是真的
因为花时间不务正业
10h 32m ago
Replied to a topic by sickoo 生活 [有一个朋友] 老丈人说结婚送房?
我反正绝对会拒绝
他有了面子有了里子,我当冤大头,我才不干
跑了下 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;
}
}
```
14h 33m ago
Replied to a topic by 0x114514 生活 这个世界太割裂了
看到学生睡觉也不管

请问要怎么管,怎么管都会被投诉,那只能不管。
直接上 x 看不就好了
x 也没那也很正常

国内 71 也没什么建党节报道,全在攻击莫名其妙的游戏
@digitv 单身当然支持
如果钱是纯粹的转移支付,那当然不支持
16h 7m ago
Replied to a topic by dsj11231 生活 哎,婚姻的意义是什么
《每个月给我转账 2500 ,算是不带孩子的补偿》 ???
给你脸了,哪来的脸要 2500
吃一堑长一智,下次就不会这么大好人了
About   ·   Help   ·   Advertise   ·   Blog   ·   API   ·   FAQ   ·   Solana   ·   1130 Online   Highest 6679   ·     Select Language
创意工作者们的社区
World is powered by solitude
VERSION: 3.9.8.5 · 13ms · UTC 18:00 · PVG 02:00 · LAX 11:00 · JFK 14:00
♥ Do have faith in what you're doing.