Add A Comment:
Just a comment, I patched P3 with the changes recommended, and still get the same segfault in a 180k atom system, though the number of GPU interaction blocks gets set to 240. I should add that even for a 22k atom system, this maximum limit is reached. If there is anything else that you think I could do, let me know, otherwise I'll just wait for the P4 release and see if that solves the issue.
The changes are quite simple, if you want to integrate them yourself. In gpu.cpp, find the function gpuBuildThreadBlockWorkList(). About half way through, you'll find a line that sets the value of gpu->sim.interaction_blocks. Immediately after that line, add the following: if (gpu->sim.interaction_blocks > 8*gpu->sim.blocks) gpu->sim.interaction_blocks = 8*gpu->sim.blocks; Next go to kFindInteractingBlocks.h. In the middle of that file (line 83 in the current code, but it might be slightly different in PR3) is an if block: if (pos < cSim.workUnits) { ... } Change it to a while loop: while (pos < cSim.workUnits) { ... pos += gridDim.x*blockDim.x; }
Thanks Peter. Would it be possible to get me a patch for this for the P3 release? I ask because I think this segfault is causing the GPU to stop responding to further OpenMM simulations, I get "launch failure" errors after the initial segfault. I'm hoping that this fix will allow me to proceed without rebooting. I understand if its complicated and I should wait for the next release.
I've checked in a fix for this, so it should work correctly in the next release.
