                 __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);

extern int ghost_cu_device;

#define CHOOSE_KERNEL(dt1,dt2) {\
    if ((SELL(mat)->T > 32) || (SELL(mat)->T == 0) || (SELL(mat)->T & (SELL(mat)->T-1)))\
        WARNING_LOG("Invalid T: %d (must be power of two <33",SELL(mat)->T);\
    if (SELL(mat)->chunkHeight == SELL(mat)->nrowsPadded) {\
        if (SELL(mat)->T > 1) {\
            INFO_LOG("ELLPACK-T kernel not available. Switching to SELL-T kernel although we have only one chunk. Performance may suffer.");\

        int j;
        v_t tmp;
        v_t *smem = (v_t *)shared;

        if (C == blockDim.x) {
            cs = chunkstart[blockIdx.x];
            tid = threadIdx.x;
