 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)->T > 128) || (SELL(mat)->T == 0) || (SELL(mat)->T & (SELL(mat)->T-1)))\
+    WARNING_LOG("Invalid T: %d (must be power of two and T <= 128)",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.");\

    }
}

template<typename m_t, typename v_t>  
__global__ void SELLT_kernel_CU_tmpl(v_t *lhs, v_t *rhs, int options, ghost_midx_t nrows, ghost_midx_t nrowspadded, ghost_midx_t *rowlen, ghost_midx_t *col, m_t *val, ghost_mnnz_t *chunkstart, ghost_midx_t *chunklen, ghost_midx_t C, int T)
{
    int i = threadIdx.x+blockIdx.x*blockDim.x;

            }\
            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) {\
