On 07/18/2013 12:57 AM, Karl Rupp wrote:
>
> In viennacl/linalg/cuda/sparse_matrix_operations.hpp:
>
> > +                unsigned int result_col_inc,
> > +                unsigned int result_row_size,
> > +                unsigned int result_col_size,
> > +                unsigned int result_internal_rows,
> > +                unsigned int result_internal_cols) {
> > +
> > +        extern __shared__ T sh_data[];
> > +        T* sh_elements = (T*)(sh_data);
> > +        unsigned int* sh_col_indices = (unsigned int*)(sh_elements + d_mat_row_size);
> > +
> > +        for (unsigned int row  = blockIdx.x; row  < result_row_size; row += gridDim.x) {
> > +
> > +          unsigned int row_start = sp_mat_row_indices[row];
> > +          unsigned int row_end = sp_mat_row_indices[row+1];
> > +
> > +          for (unsigned int x = threadIdx.x; x < (row_end - row_start); x += blockDim.x) {
>
> What if you let all threads read the same sp_mat_element[] in a burst 
> read and then let the threads operate on the dense matrix row in 
> parallel? I think that this would solve the problems with shared memory.
>
> Outline of a kernel for many columns in the dense matrix:
>
>   * Assign each work group to one row in the result matrix
>   * Each thread starts with a private float for summing the result
>     entry r.
>   * Load sp_mat_col_indices and sp_mat_elements into thread-private
>     entries j and x (this is a fast burst-read)
>   * Fetch the respective row j from the dense matrix into a
>     thread-private entry y
>   * multiply and add: r += x * y
>   * Continue with all nonzeros of that row in the sparse matrix.
>   * Write r to the result matrix
>
> This should not need any shared memory at all and give reasonably good 
> performance as long as there are at least ~32 columns in the dense matrix.
>
> —
> Reply to this email directly or view it on GitHub 
> <https://github.com/viennacl/viennacl-dev/pull/37/files#r5255773>.
>
Hi Karl,

     Job done :) https://github.com/albertzaharovits/viennacl-dev. If 
there are any more issues, please do tell !

Regards,
Albert Z

-- 
Albert Zaharovits
+40727 883 343