c - 使用 CUDA 减少矩阵行

Question

Windows 7, NVidia GeForce 425M.

我写了一个简单的 CUDA 代码来计算矩阵的行总和。矩阵具有一维表示(指向 float 的指针)。

代码的串行版本如下(如预期的那样，它有 2 循环):

void serial_rowSum (float* m, float* output, int nrow, int ncol) {
    float sum;
    for (int i = 0 ; i < nrow ; i++) {
        sum = 0;
        for (int j = 0 ; j < ncol ; j++)
            sum += m[i*ncol+j];
        output[i] = sum;
    }
}

在 CUDA 代码中，我调用内核函数逐行扫描矩阵。下面是内核调用片段:

dim3 threadsPerBlock((unsigned int) nThreadsPerBlock); // has to be multiple of 32
dim3 blocksPerGrid((unsigned int) ceil(nrow/(float) nThreadsPerBlock)); 

kernel_rowSum<<<blocksPerGrid, threadsPerBlock>>>(d_m, d_output, nrow, ncol);

以及执行行的并行求和的内核函数(仍然有 1 循环):

__global__ void kernel_rowSum(float *m, float *s, int nrow, int ncol) {

    int rowIdx = threadIdx.x + blockIdx.x * blockDim.x;

    if (rowIdx < nrow) {
        float sum=0;
        for (int k = 0 ; k < ncol ; k++)
            sum+=m[rowIdx*ncol+k];
        s[rowIdx] = sum;            
    }

}

到目前为止一切顺利。串行和并行 (CUDA) 结果相等。

重点是 CUDA 版本的计算时间几乎是串行版本的两倍，即使我更改了 nThreadsPerBlock 参数:我测试了 nThreadsPerBlock 来自 32 到 1024(我的卡允许的每个 block 的最大线程数)。

IMO，矩阵维度足够大以证明并行化:90,000 x 1,000。

下面，我使用不同的 nThreadsPerBlock 报告了串行和并行版本的运行时间。平均 100 样本的时间以 毫秒 报告:

矩阵:nrow = 90000 x ncol = 1000

序列号:每个样本平均耗用时间(以毫秒为单位)(100 样本):289.18。

CUDA(32 ThreadsPerBlock):每个样本平均耗用时间(以毫秒为单位)(100 样本):497.11。

CUDA(1024 ThreadsPerBlock):每个样本平均耗用时间(以毫秒为单位)(100 样本):699.66。

以防万一，带有32/1024 nThreadsPerBlock 的版本是最快/最慢的。

我知道从主机复制到设备时会产生开销，反之亦然，但速度慢可能是因为我没有实现最快的代码。

因为我远不是 CUDA 专家:

我是否为此任务编写了最快的版本？我怎样才能改进我的代码？我可以摆脱内核函数中的循环吗？

任何想法表示赞赏。

编辑 1

虽然我描述了一个标准的rowSum，但我对具有(0; 1} 值，如 rowAND/rowOR。也就是说，它不允许我利用 cuBLAS 乘以 1 的 COL 列 vector 技巧，正如一些评论员所建议的那样。

编辑 2

根据用户其他用户的建议并在此认可:

忘记尝试编写您自己的函数，改用 Thrust 库，魔法就来了。

Answer 1

既然你提到你需要通用的归约算法，而不是只求和。我将尝试在这里提供 3 种方法。内核方法可能具有最高的性能。推力法最容易实现。 cuBLAS 方法仅适用于 sum 并且具有良好的性能。

内核方法

Here's a very good doc介绍如何优化标准并行缩减。标准还原可分为2个阶段。

1. 多个线程 block 各减少一部分数据；
2. 一个线程 block 从阶段 1 的结果减少到最后的 1 个元素。

对于您的多重归约(减少垫子的行数)问题，只有第 1 阶段就足够了。这个想法是每个线程 block 减少 1 行。对于每个线程 block 多行或每个多个线程 block 1 行的进一步考虑，您可以引用 paper provided by @Novak .这可能会进一步提高性能，尤其是对于形状不佳的矩阵。

推力进近

一般的multi-reduction可以通过thrust::reduction_by_key来完成几分钟后。您可以在此处找到一些讨论 Determining the least element and its position in each matrix column with CUDA Thrust .

但是thrust::reduction_by_key不假设每行的长度相同，因此您会受到性能损失。另一个帖子How to normalize matrix columns in CUDA with max performance?给出 thrust::reduction_by_key 之间的分析比较和 cuBLAS 处理行总和的方法。它可以让您对性能有一个基本的了解。

cuBLAS 方法

矩阵 A 的行/列之和可以看作是矩阵 vector 乘法，其中 vector 的元素都是 1。它可以用下面的matlab代码表示。

y = A * ones(size(A,2),1);

哪里y是 A 的行的总和。

cuBLAS 库提供高性能矩阵 vector 乘法函数 cublas<t>gemv() 对于这个操作。

计时结果表明，该例程仅比简单读取一次A的所有元素慢10~50%，可以看作是该操作的理论性能上限。