52efa12fda
The token loop reads sK[] in the state update (bottom of loop) but has no barrier before the next iteration overwrites sK[] (top of loop). Without an explicit memory fence, hardware/compiler reordering can cause non-deterministic reads from shared memory. Per review: with this barrier in place, the prior __syncthreads() after the cross-warp reduction and the one immediately after loop exit are both redundant. The new barrier is a full block-level fence that also orders all_sum1/all_sum2 reads vs. the next iteration's writes, and every thread reaches it before leaving the loop. Both redundant barriers removed. No performance impact — GPU utilization is 31-33% during inference, bottlenecked by CPU MoE expert computation, not the CUDA kernel. Co-authored-by: Mark Alonzo <mark.alonzo@outlook.com> |
||
|---|---|---|
| .. | ||
| cmake | ||
| include | ||
| src | ||
| .gitignore | ||
| CMakeLists.txt | ||