P99 代码有误 #40
Description
樊老师你好，第9章原子函数的合理使用中，neighbor2gpu.cu程序find_neighbor_atomic函数内，d_NN[n1] = 0;这一行代码可能与d_NL[n2 * MN + atomicAdd(&d_NN[n2], 1)] = n1;这一行发生冲突，导致存储的d_NL有误。
假设n1线程执行过程中，修改了某个n2对应的d_NN[n2]，但若n2线程随后刚开始执行，则会在开始处执行d_NN[n2] = 0;这行代码，这导致了之前对d_NN[n2]的累加修改又重置为了零，因为这时又从0开始累加，会使得实际存储的某个粒子的邻居的数目减少，即重置为0之前的存储邻居会被新的邻居粒子覆盖。
例如在我的本地测试中，第14514个粒子的坐标为49.2969659427102 154.94173688193416，在neighbor1cpu.cu版本的程序中测试，其存储下来的neighbor.txt中对应的数据为3 14316 14514 14516 NaN NaN NaN NaN NaN NaN NaN，但在原子函数的版本中则输出为2 14514 14516 NaN NaN NaN NaN NaN NaN NaN NaN，少了14316这个邻居粒子，若删除d_NN[n1] = 0;这一行，并在每次调用核函数前，统一把d_NN置为0即能得到正确的结果：
if (atomic)
{
cudaMemset(d_NN, 0, sizeof(int) * N);
find_neighbor_atomic<<<grid_size, block_size>>>
(d_NN, d_NL, d_x, d_y, N, cutoff_square);
}
不知理解是否有误，还请指教。