         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();

            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(ipos + 1, jpos, kpos, strides)]) - __ldg(&in[index(ipos, jpos, kpos, strides)])) >
                0) {
                flx[cache_index(iblock_pos, jblock_pos)] = 0.;
            }

        __syncthreads();

        if (is_in_domain< 0, 0, 0, 0 >(iblock_pos, jblock_pos, block_size_i, block_size_j)) {
            out[index_] =
                __ldg(&in[index_]) -
                coeff[index_] *
                    (flx[cache_index(iblock_pos, jblock_pos)] - flx[cache_index(iblock_pos - 1, jblock_pos)] +
                        fly[cache_index(iblock_pos, jblock_pos)] - fly[cache_index(iblock_pos, jblock_pos - 1)]);
        }

        index_ += index(0,0,1, strides);
    }
}

