#define FPTYPE float

__kernel void
spmv_csr_vector_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 rowDelimiters,
                       const int dim, const int vecWidth, __global FPTYPE * restrict out)
{
    // Thread ID in block
    int t = get_local_id(0);
    // Thread ID within warp
    int id = t & (vecWidth-1);
    // One row per warp
    int vecsPerBlock = get_local_size(0) / vecWidth;
    int myRow = (get_group_id(0) * vecsPerBlock) + (t / vecWidth);

    __local volatile FPTYPE partialSums[128];
    partialSums[t] = 0;

    if (myRow < dim)
    {
        int vecStart = rowDelimiters[myRow];
        int vecEnd = rowDelimiters[myRow+1];
        FPTYPE mySum = 0;
        for (int j= vecStart + id; j < vecEnd;
             j+=vecWidth)
        {
            int col = cols[j];
#ifdef USE_TEXTURE
            mySum += val[j] * texFetch(vec,col);
#else
            mySum += val[j] * vec[col];
#endif
        }

        partialSums[t] = mySum;
        barrier(CLK_LOCAL_MEM_FENCE);

        // Reduce partial sums
	int bar = vecWidth / 2;
	while(bar > 0)
	{
	    if (id < bar) partialSums[t] += partialSums[t+ bar];
	    barrier(CLK_LOCAL_MEM_FENCE);
	    bar = bar / 2;
	}

        // Write result
        if (id == 0)
        {
            out[myRow] = partialSums[t];
        }
    }
}

