加速涉及转置的cuSPARSE稀疏密集乘法的方法答案

【问题标题】：method to speed up cuSPARSE sparse-dense multiplication involving transpose加速涉及转置的cuSPARSE稀疏密集乘法的方法
【发布时间】：2016-11-05 02:41:41
【问题描述】：

我使用 cusparse 和 cublas 来计算稀疏密集乘法：C = A' * B。

A 是一个 M*N 稀疏矩阵

B是一个M*S密集矩阵

M = 9,633,792，N = 617,004，nnz 为 28,901,376，S = 3

我尝试了不同的方法来让它更快，

A以CSR格式存储，使用cusparseScsrmm计算A'*B，耗时180ms
A' = At 以 CSR 格式存储，使用 cusparseScsrmm2 计算 At*(B')'，然后转置 B 以提高矩阵 B 的内存访问，根据文档，如果 op(B ) = B^T，只支持op(A) = A，所以我提前以CSR形式存储At，转置B需要8ms，计算At*(B')'需要4ms，一共12ms。
A' = At以CSR格式存储，使用cusparseScsrmm计算A'*B，耗时8ms。

A 在迭代中是常数，所以不能考虑对 A 的操作时间，但要考虑对 B 的操作时间。更具体地说，A 是一个二元矩阵，每行有 3 个非零值。

所以我在徘徊，有什么方法可以加快速度吗？ 4ms 可能是可以接受的。例如，改善矩阵 B 的内存访问但不耗时。我也考虑过用常量内存来存储A，但是cuda好像只有64K的常量内存，或者用纹理内存来存储B，不过是只读内存，可能不太合适。

/****补充***/

我用的GPU是GTX TITAN X，我用cublasSgeam转置矩阵B

【问题讨论】：

如何转置B？宝石（）？它似乎比可能的峰值慢得多。
是的，我会按照你的建议写一个内核@kangshiyin

标签： memory matrix cuda sparse-matrix

【解决方案1】：

我不知道您使用的是哪种设备。但与高端 GPU 上的单个 D2D 副本 B 相比，矩阵转置速度相当慢。

M*S*sizeof(float)/8e-3/1e9=14GB/s

如果您使用cublas_geam() 转置矩阵B，这可能意味着矩阵太薄并且例程没有针对这种情况进行很好的优化。您可以实现自己的转置内核并针对 3-col 矩阵对其进行优化。通用矩阵转置仅适用于 32 倍数或远大于 32 的维度。但它的代码是实现您自己的代码的良好开端。

https://devblogs.nvidia.com/parallelforall/efficient-matrix-transpose-cuda-cc/

【讨论】：

我的GPU是GTX TITAN X，我用cublas_geam()转置矩阵B，谢谢你的建议，我去试试
感谢您的建议，我编写了一个简单的转置核函数，大约需要 1.9 毫秒。我这里还有一个问题，cuda在Release和Debug模式下同时运行是否正常，我运行cuda程序500次，但不同模式下时间相同。我通过 'cmake -DCMAKE_BUILD_TYPE=Release' 或 'cmake -DCMAKE_BUILD_TYPE=Debug' 设置构建模式。
@zjhthu 是的，通常是一样的。