从 CPU 并行中学到的 GPU kernel 经验：依赖链、浮点非结合性、伪共享与内存受限循环
CPU parallelism lessons that carry over to GPU kernels: dependency chains, FP non-associativity, false sharing, memory-bound loops [P]
一位作者在免费 GPU kernel 工程系列第二期中，用纯 CPU 代码讲解可迁移到 GPU 的核心概念：单个累加器构成依赖链，需用多个累加器拆出独立链；浮点加法不满足结合律，编译器默认保持原顺序，GPU 归约同样存在这种非确定性。他还讨论了分支预测、Amdahl 定律、伪共享和内存受限的点积循环。
Session 2 of my free GPU kernel engineering series (my own channel, earlier post: [LINK TO FIRST POST]) covers CPU parallelism, with no GPU code. The idea is that most GPU performance ideas are CPU ideas turned up, so here are the parts I found most useful on their own.
A single accumulator is a dependency chain. In `sum += a[i] * b[i]` every add waits on the previous one, so the loop uses a fraction of the core. Independent ops needed in flight is about latency x throughput, and four accumulators give four independent chains with the same arithmetic. Measure this with an in-cache array and -fno-tree-vectorize, or you'll be timing two effects at once.
The compiler won't fix it for you. FP addition isn't associative, so it must keep your order unless you allow otherwise (-ffast-math, or `#pragma omp simd reduction(+:sum)` per loop). The same non-determinism appears in GPU reductions.
Branch cost is predictability. The sorted-vs-random loop showed a [X]x gap at -O0 on my machine ([CPU], GCC [version]) and mostly vanished at -O2, because the compiler removed the branch. Check what it generated before blaming the hardware.
Threads: Amdahl says 90% parallel on 8 cores gives about 4.7x, not 8. Per-thread partials in one small array cause false sharing, because they land on one 64-byte cache line that bounces between cores. Padding fixed it for me ([X]x faster at [N] threads), though the gap depends on whether threads sit on separate physical cores.
A dot product is memory-bound. It does 2 FLOPs per 8 bytes, so at [BW] GB/s it caps near [Y] GFLOP/s no matter how much SIMD or threading you add. Hiding memory latency needs latency x bandwidth bytes in flight, which is the GPU's reason for running thousands of threads.
Video: https://youtu.be/_xE1_2FFTC0
Caveats: numbers are from one machine, intrinsics assume AVX2, and the latency/throughput figures are typical Intel values.
Is anything wrong or oversimplified, especially the false-sharing explanation? And which CPU lessons transferred best to GPU work for you?
submitted by /u/Winter_Mistake_3185
[link] [留言]
来源：r/MachineLearning · reddit.com