Skip to content

GitLab

Commit 2d7fb355 authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed error building neighbor list with triclinic box

parent 509fa77e
Hide whitespace changes
Inline Side-by-side
Showing with 3 additions and 2 deletions
platforms/cuda/src/kernels/findInteractingBlocks.cu
View file @ 2d7fb355
...@@ -243,6 +243,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea ...@@ -243,6 +243,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
for (int block2Base = block1+1; block2Base < NUM_BLOCKS; block2Base += 32) { for (int block2Base = block1+1; block2Base < NUM_BLOCKS; block2Base += 32) {
int block2 = block2Base+indexInWarp; int block2 = block2Base+indexInWarp;
bool includeBlock2 = (block2 < NUM_BLOCKS); bool includeBlock2 = (block2 < NUM_BLOCKS);
bool forceInclude = false;
if (includeBlock2) { if (includeBlock2) {
real4 blockCenterY = sortedBlockCenter[block2]; real4 blockCenterY = sortedBlockCenter[block2];
real4 blockSizeY = sortedBlockBoundingBox[block2]; real4 blockSizeY = sortedBlockBoundingBox[block2];
...@@ -260,7 +261,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea ...@@ -260,7 +261,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
// If there's any possibility we might have missed it, do a detailed check. // If there's any possibility we might have missed it, do a detailed check.
if (periodicBoxSize.z/2-blockSizeX.z-blockSizeY.z < PADDED_CUTOFF || periodicBoxSize.y/2-blockSizeX.y-blockSizeY.y < PADDED_CUTOFF) if (periodicBoxSize.z/2-blockSizeX.z-blockSizeY.z < PADDED_CUTOFF || periodicBoxSize.y/2-blockSizeX.y-blockSizeY.y < PADDED_CUTOFF)
includeBlock2 = true; includeBlock2 = forceInclude = true;
#endif #endif
if (includeBlock2) { if (includeBlock2) {
unsigned short y = (unsigned short) sortedBlocks[block2].y; unsigned short y = (unsigned short) sortedBlocks[block2].y;
...@@ -291,7 +292,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea ...@@ -291,7 +292,7 @@ extern "C" __global__ void findBlocksWithInteractions(real4 periodicBoxSize, rea
#ifdef USE_PERIODIC #ifdef USE_PERIODIC
APPLY_PERIODIC_TO_DELTA(atomDelta) APPLY_PERIODIC_TO_DELTA(atomDelta)
#endif #endif
int atomFlags = BALLOT(atomDelta.x*atomDelta.x+atomDelta.y*atomDelta.y+atomDelta.z*atomDelta.z < (PADDED_CUTOFF+blockCenterY.w)*(PADDED_CUTOFF+blockCenterY.w)); int atomFlags = BALLOT(forceInclude || atomDelta.x*atomDelta.x+atomDelta.y*atomDelta.y+atomDelta.z*atomDelta.z < (PADDED_CUTOFF+blockCenterY.w)*(PADDED_CUTOFF+blockCenterY.w));
int interacts = 0; int interacts = 0;
if (atom2 < NUM_ATOMS && atomFlags != 0) { if (atom2 < NUM_ATOMS && atomFlags != 0) {
int first = __ffs(atomFlags)-1; int first = __ffs(atomFlags)-1;
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment