+#define CUSP_USE_TEXTURE_MEMORY
+
 #include <cusp/dia_matrix.h>
 #include <cusp/csr_matrix.h>
+#include <cusp/io/matrix_market.h>
 
 #include <thrust/sequence.h>
 #include <thrust/fill.h>

        unbind_x(x);
}
 
template <typename IndexType, typename ValueType, unsigned int ThreadsPerVector>
void benchmark_csr_vector()
{
    printf("%4d,", (int) ThreadsPerVector);

    cusp::array1d<ValueType, cusp::device_memory> x(N);
    cusp::array1d<ValueType, cusp::device_memory> y(N);

    for(IndexType D = 1; D <= max_diagonals; D++)
    {
        // create banded matrix with D diagonals
        const IndexType NNZ = N * D - (D * (D - 1)) / 2;
        cusp::dia_matrix<IndexType, ValueType, cusp::host_memory> dia(N, N, NNZ, D, N);
        thrust::sequence(dia.diagonal_offsets.begin(), dia.diagonal_offsets.end());
        thrust::fill(dia.values.values.begin(), dia.values.values.end(), 1);

        // convert to CSR
        cusp::csr_matrix<IndexType, ValueType, cusp::device_memory> csr(dia);
    
        // time several SpMV iterations
        timer t;
        for(size_t i = 0; i < num_iterations; i++)
            perform_spmv<true, ThreadsPerVector>(csr, thrust::raw_pointer_cast(&x[0]), thrust::raw_pointer_cast(&y[0]));
        cudaThreadSynchronize();

        float sec_per_iteration = t.seconds_elapsed() / num_iterations;
        float gflops = (NNZ/sec_per_iteration) / 1e9;

        printf("%4.1f,", gflops);
    }

    std::cout << std::endl;
}


int main(int argc, char** argv)
{
    printf("# of diagonals vary along rows, # of threads per vector varies along column\n");
    printf("    ,");
    for(size_t D = 1; D <= max_diagonals; D++)
        printf("%4d,", (int) D);
    std::cout << std::endl;

    benchmark_csr_vector<int, float,  2>();
    benchmark_csr_vector<int, float,  4>();
    benchmark_csr_vector<int, float,  8>();
    benchmark_csr_vector<int, float, 16>();
    benchmark_csr_vector<int, float, 32>();

    return 0;
}


#include <cusp/detail/device/spmv.h>
#include "../timer.h"

template <bool UseCache, unsigned int THREADS_PER_VECTOR, typename IndexType, typename ValueType>
void perform_spmv(const cusp::csr_matrix<IndexType,ValueType,cusp::device_memory>& csr, 
