                 __syncthreads();
             }
         }
-        
+
         if (threadIdx.y == 0) {
             if (options & GHOST_SPMVM_AXPY)
                 lhs[i] = axpy<v_t,float>(lhs[i],axpy<v_t,float>(smem[tib],smem[tib+1],1.f),1.f);

            }\
            dim3 block(SELL_CUDA_THREADSPERBLOCK/SELL(mat)->T,SELL(mat)->T);\
            SELLT_kernel_CU_tmpl\
                <dt1,dt2>\
                <<< SELL_CUDA_NBLOCKS,block,reqSmem >>>\
                ((dt2 *)lhs->CU_val,(dt2 *)rhs->CU_val,options,SELL(mat)->cumat->nrows,SELL(mat)->cumat->nrowsPadded,SELL(mat)->cumat->rowLenPadded,SELL(mat)->cumat->col,(dt1 *)SELL(mat)->cumat->val,SELL(mat)->cumat->chunkStart,SELL(mat)->cumat->chunkLen,SELL(mat)->chunkHeight,SELL(mat)->T);\
        } else {\
            SELL_kernel_CU_tmpl\
                <dt1,dt2>\
                <<< SELL_CUDA_NBLOCKS,SELL_CUDA_THREADSPERBLOCK >>>\
                ((dt2 *)lhs->CU_val,(dt2 *)rhs->CU_val,options,SELL(mat)->cumat->nrows,SELL(mat)->cumat->nrowsPadded,SELL(mat)->cumat->rowLen,SELL(mat)->cumat->col,(dt1 *)SELL(mat)->cumat->val,SELL(mat)->cumat->chunkStart,SELL(mat)->cumat->chunkLen,SELL(mat)->chunkHeight,SELL(mat)->T);\
        }\
    }\
}
 
template<typename m_t, typename v_t>  
__global__ void SELL_kernel_CU_ELLPACK_tmpl(v_t *lhs, v_t *rhs, int options, int nrows, int nrowspadded, ghost_midx_t *rowlen, ghost_midx_t *col, m_t *val, ghost_mnnz_t *chunkstart, ghost_midx_t *chunklen, int C, int T)
{
    UNUSED(C);

            }\
            dim3 block(SELL_CUDA_THREADSPERBLOCK/SELL(mat)->T,SELL(mat)->T);\
            SELLT_kernel_CU_tmpl\
            <dt1,dt2>\
            <<< SELL_CUDA_NBLOCKS,block,reqSmem >>>\
            ((dt2 *)lhs->CU_val,(dt2 *)rhs->CU_val,options,SELL(mat)->cumat->nrows,SELL(mat)->cumat->nrowsPadded,SELL(mat)->cumat->rowLenPadded,SELL(mat)->cumat->col,(dt1 *)SELL(mat)->cumat->val,SELL(mat)->cumat->chunkStart,SELL(mat)->cumat->chunkLen,SELL(mat)->chunkHeight,SELL(mat)->T);\
        } else {\
            SELL_kernel_CU_ELLPACK_tmpl\
            <dt1,dt2>\
            <<< SELL_CUDA_NBLOCKS,SELL_CUDA_THREADSPERBLOCK >>>\
            ((dt2 *)lhs->CU_val,(dt2 *)rhs->CU_val,options,SELL(mat)->cumat->nrows,SELL(mat)->cumat->nrowsPadded,SELL(mat)->cumat->rowLen,SELL(mat)->cumat->col,(dt1 *)SELL(mat)->cumat->val,SELL(mat)->cumat->chunkStart,SELL(mat)->cumat->chunkLen,SELL(mat)->chunkHeight,SELL(mat)->T);\
        }\
    }else{\
        if (SELL(mat)->T > 1) {\
