#define FPTYPE float

__kernel void
spmv_ellpackr_kernel(__global const FPTYPE * restrict val,
#ifdef USE_TEXTURE
                     image2d_t vec,
#else
                     __global const  FPTYPE * restrict vec,
#endif
                     __global const int * restrict cols,
                     __global const int * restrict rowLengths,
                     const int dim, __global FPTYPE * restrict out)
{
    int t = get_global_id(0);

    if (t < dim)
    {
        FPTYPE result = 0.0;
        int max = rowLengths[t];
        for (int i = 0; i < max; i++)
        {
            int ind = i * dim + t;
#ifdef USE_TEXTURE
	          result += val[ind] * texFetch(vec,cols[ind]);
#else
	          result += val[ind] * vec[cols[ind]];
#endif
        }
        out[t] = result;
    }
}



