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

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