关于P102 find_neighbor_no_atomic 中 d_NL 访问模式的疑问 #49
Description
老师您好，
在书的示例代码 find_neighbor_no_atomic 核中，有如下一行代码：
d_NL[(count++) * N + n1] = n2;书中对此的解释是：我们改变了邻居列表中的数据排列方式：将d_NL[n1*MN+count++]改成了d_NL[(count++) * N + n1] 。因为n1的变化步调与threadIdx.x一致，这样修改以后，对全局内存数组d_NL的访问将是合并的（不考虑数据不对齐带来的影响）
我对这个解释有些困惑。根据我的理解，内存合并访问要求同一个warp内的线程访问连续的全局内存地址。
对于 d_NL[(count++) * N + n1] 这个访问模式：线程的索引 n1 确实与 threadIdx.x 相关。但是，内存地址的计算方式是 (count * N + n1)。由于 count 的值取决于每个粒子 n1 实际找到的邻居数量，对于同一个warp内的不同线程，它们的 count 值在循环过程中是不同步且不可预测的。这导致同一个warp内的线程访问的内存地址 ((count_A * N + n1_A), (count_B * N + n1_B), ...) 是完全不连续且跳跃的，步长为 N 的倍数，这似乎是非合并的访问模式。
因此，我怀疑书中的这段解释可能存在笔误或需要更详细的说明。是否是我对“合并访问”的理解有偏差，还是这里的解释确实有误？
非常感谢您和这本书带给我的启发，期待能有机会得到您的解答！