         jblock_pos = threadIdx.x / PADDED_BOUNDARY;
     }
 
+    int index_ = index(ipos, jpos, 0, strides);
+
 // 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

        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)]));
        }

        __syncthreads();

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