Question on design choice in merge-spmv shared memory loading

[darius513]

Great work!

I have a few questions regarding the merge-spmv implementation. In particular, I noticed that in the csrgemv_merge function, lines 366–371 contain the following code:

#pragma unroll
for (int j = 0; j < TASKS_PER_THREAD; j++) {
    sa[j*THREADS_PER_BLOCK + threadIdx.x] =
        alpha * a[kmin + j*THREADS_PER_BLOCK + threadIdx.x] *
        x[colidx[kmin + j*THREADS_PER_BLOCK + threadIdx.x]];
    srowptr[j*THREADS_PER_BLOCK + threadIdx.x] =
        rowptr[imin + j*THREADS_PER_BLOCK + threadIdx.x];
}

When loading the row offsets (list A) and the matrix–vector dot products (list B) into shared memory, this code appears to perform additional work. Starting from the initial point ((imin, kmin)), it advances to the right and downward by THREADS_PER_BLOCK * TASKS_PER_THREAD, meaning that list A and list B each consume THREADS_PER_BLOCK * TASKS_PER_THREAD elements.

This behavior seems somewhat different from the implementation described in the original merge-spmv paper
(https://ieeexplore.ieee.org/document/7877136). In merge-spmv, after computing the coordinates of the start and end points, the total number of consumed elements from list A and list B is THREADS_PER_BLOCK * TASKS_PER_THREAD.

Could you please share the rationale behind this design choice? Do you expect these two approaches to exhibit any performance differences in practice?

[jamtrott]

Hi @darius513, I believe you are right in that the current code will often prefetch more data to shared memory than is strictly needed. If I recall correctly, we fetch enough data for the row pointers/offsets (list A) to cover a worst case where every row is empty. In practice, a much smaller number of row pointers will be needed.

An earlier version of the code did follow the procedure as in the original paper, which you outlined above. If I remember correctly, this requires additional branching or logic during the prefetching to shared memory (e.g., to decide how many entries to fetch). In the end, it was simply faster to prefetch a fixed amount and let the compiler unroll the loop automatically.

Unfortunately, I do not have any specific measurements at hand anymore. Please share your findings, if you decide to investigate this more closely.

Thank you for the question!