Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

P99 代码有误 #40

Open
yuchangminghit opened this issue Feb 6, 2025 · 2 comments
Open

P99 代码有误 #40

yuchangminghit opened this issue Feb 6, 2025 · 2 comments

Comments

@yuchangminghit
Copy link

樊老师你好,第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);
}

不知理解是否有误,还请指教。

@fever-Wong
Copy link

fever-Wong commented Feb 6, 2025 via email

@brucefan1983
Copy link
Owner

谢谢指出。你的理解是正确的,我的书稿和代码确实弄错了。我会在主页列上这个刊误。

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

3 participants