title: Sparse matrix-vector multiplication in CUDA
body: <p>I'm trying to implement a matrix-vector Multiplication on GPU (using CUDA).  </p>
  
  <p>In my C++ code (CPU), I load the matrix as a dense matrix, and then I perform the matrix-vector multiplication using CUDA. I'm also using shared memory to improve the performance.  </p>
  
  <ol>
  <li>How can I load the matrix in an efficient way, knowing that my matrix is a sparse matrix?</li>
  </ol>
  
  <p>Below is my C++ function to load the matrix:  </p>
  
  <pre><code>int readMatrix( char* filename, float* &amp;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 &gt;&gt; h &gt;&gt; w &gt;&gt; num_entries;
          cout &lt;&lt; w &lt;&lt; " " &lt;&lt; h &lt;&lt; " " &lt;&lt; num_entries &lt;&lt; "\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 &lt; num_entries; i++ ){
  
              if( file.eof() ) break;
  
              file &gt;&gt; y &gt;&gt; x &gt;&gt; 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 &lt;&lt; "\nFile read successfully\n"; 
          else
              std::cout &lt;&lt; "\nFile read successfully but seems defective:\n num entries read = " &lt;&lt; i &lt;&lt; ", entries epected = " &lt;&lt; num_entries &lt;&lt; "\n"; 
  
          // print first few elements
          if( w == h ){
              for( unsigned int i = 0; i &lt; w; i++ ){
  
                  printf("\n");
                  for( unsigned int j = 0; j &lt; h; j++ ){
  
                      printf("%.2f ", matrix[ j + w * i ] );
                  }
              }   
          }
          else{   
  
              printf("\n");
              for( unsigned int j = 0; j &lt; h; j++ ){
  
                  printf("%.2f ", matrix[ j ] );
              }
          }
  
      } else {
  
          std::cout &lt;&lt; "Unable to open file\n";
          return false;
      }
  
      return true;
  }
  </code></pre>
  
  <p>Below is my CUDA Kernel function that handles the matrix-vector multiplication:   </p>
  
  <pre><code>__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 &lt; (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 &lt; dim) {
          temp = 0.0;
          //dot product (multiplication)
          for (int i = 0; i &lt; dim; i++){
              temp += A[idx * dim + i] * vec[i];
          }
           x[idx] = temp;
      } 
  
  }
  </code></pre>
  
  <ol start="2">
  <li>What are the necessary changes that I have to make on my CUDA code to take into account that my matrix is a sparse matrix?  </li>
  <li>I found out from a forum that we can also use padding to be able to optimize the performance, but this requires me to change the way I read the matrix / sort the matrix. Any ideas how to implement this padding in the way I read the matrix and perform the calculation?  </li>
  </ol>