Avoid potential integer overflow in warp arithmetic 31/4631/1
authorJim Phillips <jim@ks.uiuc.edu>
Fri, 28 Sep 2018 20:44:57 +0000 (15:44 -0500)
committerJim Phillips <jim@ks.uiuc.edu>
Fri, 28 Sep 2018 20:44:57 +0000 (15:44 -0500)
Replace blockDim.x*foo/WARPSIZE with blockDim.x/WARPSIZE*foo
since we know blockDim.x is a multiple of WARPSIZE.  This avoids
overflow of blockDim.x*foo for large foo (blockIdx.x or gridDim.x).
This fix is cautionary as no bug has actually been observed.

Change-Id: I276feb01323d2b677a0456254d026681cd39f56c

src/CudaComputeGBISKernel.cu
src/CudaComputeNonbondedKernel.cu
src/CudaTileListKernel.cu

index 57ad84d..3812749 100644 (file)
@@ -315,7 +315,7 @@ GBIS_Kernel(const int numTileLists,
   TileListVirialEnergy* __restrict__ virialEnergy) {
 
   // Single warp takes care of one list of tiles
-  for (int itileList = (threadIdx.x + blockDim.x*blockIdx.x)/WARPSIZE;itileList < numTileLists;itileList += blockDim.x*gridDim.x/WARPSIZE)
+  for (int itileList = threadIdx.x/WARPSIZE + blockDim.x/WARPSIZE*blockIdx.x;itileList < numTileLists;itileList += blockDim.x/WARPSIZE*gridDim.x)
   {
     TileList tmp = tileLists[itileList];
     int iatomStart = tmp.iatomStart;
index 7c6d703..4d564f7 100644 (file)
@@ -191,7 +191,7 @@ nonbondedForceKernel(const int start, const int numTileLists,
 
   // Single warp takes care of one list of tiles
   // for (int itileList = (threadIdx.x + blockDim.x*blockIdx.x)/WARPSIZE;itileList < numTileLists;itileList += blockDim.x*gridDim.x/WARPSIZE)
-  int itileList = start + (threadIdx.x + blockDim.x*blockIdx.x)/WARPSIZE;
+  int itileList = start + threadIdx.x/WARPSIZE + blockDim.x/WARPSIZE*blockIdx.x;
   if (itileList < numTileLists)
   {
 
index b68f02f..810f348 100644 (file)
@@ -191,8 +191,7 @@ __global__ void buildSortKeys(const int numTileListsDst, const int maxTileListLe
 __global__ void fillSortKeys(const int numComputes, const int maxTileListLen,
   const int2* __restrict__ minmaxListLen, unsigned int* __restrict__ sortKeys) {
 
-  int i = (threadIdx.x + blockDim.x*blockIdx.x)/WARPSIZE;
-  if (i < numComputes) {
+  for (int i = threadIdx.x/WARPSIZE + blockDim.x/WARPSIZE*blockIdx.x;i < numComputes;i+=blockDim.x/WARPSIZE*gridDim.x) {
     const int wid = threadIdx.x % WARPSIZE;
     int2 minmax = minmaxListLen[i];
     int minlen = minmax.x;
@@ -541,7 +540,7 @@ repackTileListsKernel(const int numTileLists, const int begin_bit, const int* __
   const int wid = threadIdx.x % WARPSIZE;
 
   // One warp does one tile list
-  for (int i = (threadIdx.x + blockDim.x*blockIdx.x)/WARPSIZE;i < numTileLists;i+=blockDim.x*gridDim.x/WARPSIZE
+  for (int i = threadIdx.x/WARPSIZE + blockDim.x/WARPSIZE*blockIdx.x;i < numTileLists;i+=blockDim.x/WARPSIZE*gridDim.x
   {
     int j = tileListOrder[i];
     int start = tileListPos[i];