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

            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)]) >
                0) {
                flx[cache_index(iblock_pos, jblock_pos)] = 0.;
            }


        if (is_in_domain< 0, 0, 0, 0 >(iblock_pos, jblock_pos, block_size_i, block_size_j)) {
            out[index(ipos, jpos, kpos, strides)] =
                __ldg(&in[index(ipos, jpos, kpos, strides)]) -
                coeff[index(ipos, jpos, kpos, strides)] *
                    (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)]);
