     int kH, int kW, int tid, int nThreads) {
 
   // Copy kernel in global memory to shared memory
-  int nTotal = kW * kH;
-  int avgToLoad = (kW * kH) / nThreads + 1;
-
-  for (int i=0; i<avgToLoad; ++i) {
-    int id = tid + i * nThreads;
-
-    if (id >= nTotal) break;
-
-    int xx = id / kH,
-	yy = id % kH;
+  for (; tid<kW * kH; tid += nThreads) {
+    int xx = tid / kH,
+	yy = tid % kH;
 
     if (xx >= kW || yy >= kH) continue;
 

  grids.z = nOutputs;

  size_t SHM_SIZE = getSuitableShmConfig(grids, threads, imgOut.m, imgOut.n);

  for (size_t i=0; i<nInputs; ++i)
    for (size_t b=0; b<batch_size; ++b) {

  downsample_kernel<<<grids, threads>>>(
      output.getData(),
      x.getData(),
      (uint8_t) scale, (uint8_t) H, (uint8_t) W);

  CCE(cudaDeviceSynchronize());

