             tid = threadIdx.x%C;
         }
         int j;
-        v_t tmp;
-        zero<v_t>(tmp);
+        v_t *tmp = ((v_t *)shared)+ncols*threadIdx.x;
+        m_t matval;
+        
+        int c;
+        for (c=0; c<ncols; c++) {
+            zero<v_t>(tmp[c]);
+        }
 
         for (j=0; j<rowlen[i]; j++) {
-            tmp = axpy<v_t,m_t>(tmp, rhs[col[cs + tid + j*C]], val[cs + tid + j*C]);
-        }
-
-        if (!do_shift) {
-            if (!do_scale) {
-                if (do_axpy) {
-                    lhs[i] = axpy<v_t,float>(lhs[i],tmp,1.f);
-                } else if (do_axpby) {
-                    lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),tmp,1.f);
-                } else {
-                    lhs[i] = tmp;
+            matval = val[cs+tid+j*C];
+            for (c=0; c<ncols; c++) {
+                tmp[c] = axpy<v_t,m_t>(tmp[c], rhs[rhs_lda*c+col[cs + tid + j*C]], matval);
+            }
+        }
+
+        for (c=0; c<ncols; c++) {
+            if (!do_shift) {
+                if (!do_scale) {
+                    if (do_axpy) {
+                        lhs[lhs_lda*c+i] = axpy<v_t,float>(lhs[lhs_lda*c+i],tmp[c],1.f);
+                    } else if (do_axpby) {
+                        lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),tmp[c],1.f);
+                    } else {
+                        lhs[i] = tmp[c];
+                    }
+                } else {
+                    if (do_axpy) {
+                        lhs[i] = axpy<v_t,float>(lhs[i],scale<v_t>(alpha,tmp[c]),1.f);
+                    } else if (do_axpby) {
+                        lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),scale<v_t>(alpha,tmp[c]),1.f);
+                    } else {
+                        lhs[i] = scale<v_t>(alpha,tmp[c]);
+                    }
                 }
             } else {
-                if (do_axpy) {
-                    lhs[i] = axpy<v_t,float>(lhs[i],scale<v_t>(alpha,tmp),1.f);
-                } else if (do_axpby) {
-                    lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),scale<v_t>(alpha,tmp),1.f);
-                } else {
-                    lhs[i] = scale<v_t>(alpha,tmp);
-                }
-            }
-        } else {
-            if (!do_scale) {
-                if (do_axpy) {
-                    lhs[i] = axpy<v_t,float>(lhs[i],axpy<v_t,v_t>(tmp,rhs[i],scale2<v_t,float>(shift,-1.f)),1.f);
-                } else if (do_axpby) {
-                    lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),axpy<v_t,v_t>(tmp,rhs[i],scale2<v_t,float>(shift,-1.f)),1.f);
-                } else {
-                    lhs[i] = axpy<v_t,v_t>(tmp,rhs[i],scale2<v_t,float>(shift,-1.f));
-                }
-            } else {
-                if (do_axpy) {
-                    lhs[i] = axpy<v_t,float>(lhs[i],scale<v_t>(alpha,axpy<v_t,v_t>(tmp,rhs[i],scale2<v_t,float>(shift,-1.f))),1.f);
-                } else if (do_axpby) {
-                    lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),scale<v_t>(alpha,axpy<v_t,v_t>(tmp,rhs[i],scale2<v_t,float>(shift,-1.f))),1.f);
-                } else {
-                    lhs[i] = scale<v_t>(alpha,axpy<v_t,v_t>(tmp,rhs[i],scale2<v_t,float>(shift,-1.f)));
-                }
-            }
-            
+                if (!do_scale) {
+                    if (do_axpy) {
+                        lhs[i] = axpy<v_t,float>(lhs[i],axpy<v_t,v_t>(tmp[c],rhs[i],scale2<v_t,float>(shift,-1.f)),1.f);
+                    } else if (do_axpby) {
+                        lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),axpy<v_t,v_t>(tmp[c],rhs[i],scale2<v_t,float>(shift,-1.f)),1.f);
+                    } else {
+                        lhs[i] = axpy<v_t,v_t>(tmp[c],rhs[i],scale2<v_t,float>(shift,-1.f));
+                    }
+                } else {
+                    if (do_axpy) {
+                        lhs[i] = axpy<v_t,float>(lhs[i],scale<v_t>(alpha,axpy<v_t,v_t>(tmp[c],rhs[i],scale2<v_t,float>(shift,-1.f))),1.f);
+                    } else if (do_axpby) {
+                        lhs[i] = axpy<v_t,float>(scale<v_t>(lhs[i],beta),scale<v_t>(alpha,axpy<v_t,v_t>(tmp[c],rhs[i],scale2<v_t,float>(shift,-1.f))),1.f);
+                    } else {
+                        lhs[i] = scale<v_t>(alpha,axpy<v_t,v_t>(tmp[c],rhs[i],scale2<v_t,float>(shift,-1.f)));
+                    }
+                }
+                
+            }
         }
 
     }

}

    template<typename m_t, typename v_t>  
__global__ void SELL_kernel_CU_tmpl(v_t *lhs, v_t *rhs, ghost_spmv_flags_t flags, int nrows, int nrowspadded, ghost_idx_t *rowlen, ghost_idx_t *col, m_t *val, ghost_nnz_t *chunkstart, ghost_idx_t *chunklen, int C, int T, v_t shift, v_t alpha, v_t beta, v_t *localdot, const bool do_axpy, const bool do_axpby, const bool do_scale, const bool do_shift, const bool do_localdot)
{
    UNUSED(T);
    int i = threadIdx.x+blockIdx.x*blockDim.x;

            SELLT_kernel_CU_tmpl<dt1,dt2><<<SELL_CUDA_NBLOCKS,block,reqSmem>>>((dt2 *)lhs->cu_val,(dt2 *)rhs->cu_val,flags,mat->nrows,mat->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,shift,scale,beta,(dt2 *)cu_localdot,flags&GHOST_SPMV_AXPY,flags&GHOST_SPMV_AXPBY,flags&GHOST_SPMV_SCALE,flags&GHOST_SPMV_SHIFT,flags&GHOST_SPMV_DOT);\
            /*SWITCH_BOOLS(SELLT_kernel_CU_tmpl,dt1,dt2,SELL_CUDA_NBLOCKS,block,reqSmem)*/\
        } else {\
            SELL_kernel_CU_tmpl<dt1,dt2><<<SELL_CUDA_NBLOCKS,SELL_CUDA_THREADSPERBLOCK,rhs->traits.ncols*rhs->elSize*SELL_CUDA_THREADSPERBLOCK>>>((dt2 *)lhs->cu_val,lhs->traits.nrowspadded,(dt2 *)rhs->cu_val,rhs->traits.nrowspadded,flags,mat->nrows,mat->nrowsPadded,rhs->traits.ncols,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,shift,scale,beta,(dt2 *)cu_localdot,flags&GHOST_SPMV_AXPY,flags&GHOST_SPMV_AXPBY,flags&GHOST_SPMV_SCALE,flags&GHOST_SPMV_SHIFT,flags&GHOST_SPMV_DOT);\
            /*SWITCH_BOOLS(SELL_kernel_CU_tmpl,dt1,dt2,SELL_CUDA_NBLOCKS,SELL_CUDA_THREADSPERBLOCK)*/\
        }\
    }\
