Разреженное умножение матрицы на вектор в CUDA

Я пытаюсь реализовать умножение матрицы на вектор на GPU (используя CUDA).

В моем коде C ++ (CPU) я загружаю матрицу как плотную матрицу, а затем выполняю умножение матрицы на вектор с помощью CUDA. Я также использую общую память для повышения производительности.

1. Как я могу эффективно загрузить матрицу, зная, что моя матрица является разреженной?

Ниже приведена моя функция C ++ для загрузки матрицы:

int readMatrix( char* filename, float* &matrix, unsigned int *dim = NULL, int majority = ROW_MAJOR )
{
    unsigned int w, h, x, y, num_entries;

    float val;

    std::ifstream file( filename );

    if ( file )
    {
        file >> h >> w >> num_entries;
        cout << w << " " << h << " " << num_entries << "\n";

        assert( w == h || w == 1 || h == 1 );

        if( dim != NULL ) *dim = std::max( w, h );

        matrix = new float[ w * h ];

        unsigned int i;
        for( i = 0; i < num_entries; i++ ){

            if( file.eof() ) break;

            file >> y >> x >> val;

            if( majority == ROW_MAJOR ){

                matrix[ w * y + x ] = val;

            } else if( majority == COLUMN_MAJOR ){

                matrix[ h * x + y ] = val;
            }
        }
        file.close();

        if( i == num_entries )
            std::cout << "\nFile read successfully\n"; 
        else
            std::cout << "\nFile read successfully but seems defective:\n num entries read = " << i << ", entries epected = " << num_entries << "\n"; 

        // print first few elements
        if( w == h ){
            for( unsigned int i = 0; i < w; i++ ){

                printf("\n");
                for( unsigned int j = 0; j < h; j++ ){

                    printf("%.2f ", matrix[ j + w * i ] );
                }
            }   
        }
        else{   

            printf("\n");
            for( unsigned int j = 0; j < h; j++ ){

                printf("%.2f ", matrix[ j ] );
            }
        }

    } else {

        std::cout << "Unable to open file\n";
        return false;
    }

    return true;
}

Ниже приведена моя функция ядра CUDA, которая обрабатывает умножение матрицы на вектор:

__global__ void
_cl_matrix_vector_( float *A, float *b, float *x, int dim )
{
    extern __shared__ float vec[];
    unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;
    float temp = 0.0;
    int vOffs = 0;

    //load vector into shared memory
    for (int i = 0; i < (dim/blockDim.x) + 1 ; ++i, vOffs+= blockDim.x) {
        vec[vOffs + threadIdx.x] = b[vOffs + threadIdx.x];
    }

    //make sure all threads are synchronized
     __syncthreads();

    if (idx < dim) {
        temp = 0.0;
        //dot product (multiplication)
        for (int i = 0; i < dim; i++){
            temp += A[idx * dim + i] * vec[i];
        }
         x[idx] = temp;
    } 

}

2. Какие необходимые изменения я должен внести в свой код CUDA, чтобы учесть, что моя матрица является разреженной матрицей?
3. Я узнал на форуме, что мы также можем использовать заполнение, чтобы иметь возможность оптимизировать производительность, но для этого мне нужно изменить способ чтения матрицы / сортировки матрицы. Есть идеи, как реализовать это заполнение так, как я читаю матрицу и выполняю расчет?