         if (is_in_domain< -1, 1, -1, 1 >(iblock_pos, jblock_pos, block_size_i, block_size_j)) {
 
             lap[cache_index(iblock_pos, jblock_pos)] =
-                (Real)4 * __ldg(& in[index(ipos, jpos, kpos, strides)] ) -
-                ( __ldg(& in[index(ipos + 1, jpos, kpos, strides)] ) + __ldg(& in[index(ipos - 1, jpos, kpos, strides)] ) +
-                    __ldg(&in[index(ipos, jpos + 1, kpos, strides)]) + __ldg(&in[index(ipos, jpos - 1, kpos, strides)]));
+                (Real)4 * __ldg(& in[index_] ) -
+                ( __ldg(& in[index_+index(1, 0,0, strides)] ) + __ldg(& in[index_ - index(1, 0,0, strides)] ) +
+                    __ldg(&in[index_+index(0, 1, 0, strides)]) + __ldg(&in[index_ - index(0, - 1, 0, strides)]));
         }
 
         __syncthreads();

        jblock_pos = threadIdx.x / PADDED_BOUNDARY;
    }

// flx and fly can be defined with smaller cache sizes, however in order to reuse the same cache_index function, I
// defined them here
// with same size. shared memory pressure should not be too high nevertheless

            flx[cache_index(iblock_pos, jblock_pos)] =
                lap[cache_index(iblock_pos + 1, jblock_pos)] - lap[cache_index(iblock_pos, jblock_pos)];
            if (flx[cache_index(iblock_pos, jblock_pos)] *
                    (__ldg(&in[index_+index(1, 0, 0, strides)]) - __ldg(&in[index_])) >
                0) {
                flx[cache_index(iblock_pos, jblock_pos)] = 0.;
            }
