                     }\
                 }\
             } else {\
-                if (flags & GHOST_SPMV_VSHIFT) {\
+                if (flags & (GHOST_SPMV_VSHIFT | GHOST_SPMV_SHIFT)) {\
                     if (flags & GHOST_SPMV_DOT) {\
                         CALL(func,dt1,dt2,true,false,true,true,__VA_ARGS__)\
                     } else {\

        dot3 = blockReduceSum(dot3);

        if (threadIdx.x==0) {
            localdot[0*ncols*blockDim.y*gridDim.x + col*gridDim.x + blockIdx.x] = dot1;
            localdot[1*ncols*blockDim.y*gridDim.x + col*gridDim.x + blockIdx.x] = dot2;
            localdot[2*ncols*blockDim.y*gridDim.x + col*gridDim.x + blockIdx.x] = dot3;
        }
    }
#endif

#define PROCESS_LOCALDOT(dt2_host)\
        GHOST_INSTR_START(spmv_cuda_dot_reduction)\
        int block, col;\
        INFO_LOG("Experimental local dot product with final reduction over %d blocks!",grid.x);\
        dt2_host *localdot_blocks;\
        GHOST_CALL_RETURN(ghost_malloc((void **)&localdot_blocks,sizeof(dt2_host)*rhs->traits.ncols*3*grid.x));\
        GHOST_CALL_RETURN(ghost_cu_download(localdot_blocks,cu_localdot,sizeof(dt2_host)*rhs->traits.ncols*3*grid.x));\
        _Pragma("omp parallel for private(block)")\
        for (col=0; col<rhs->traits.ncols; col++) {\
            localdot[col                      ] = 0;\
            localdot[col + 1*rhs->traits.ncols] = 0;\
            localdot[col + 2*rhs->traits.ncols] = 0;\
            for (block=0; block<grid.x; block++) {\
                localdot[col                      ] += localdot_blocks[                             col*grid.x + block];\
                localdot[col + 1*rhs->traits.ncols] += localdot_blocks[1*grid.x*rhs->traits.ncols + col*grid.x + block];\
                localdot[col + 2*rhs->traits.ncols] += localdot_blocks[2*grid.x*rhs->traits.ncols + col*grid.x + block];\
            }\
        }\
        free(localdot_blocks);\
