使用CUDA减少矩阵行

Question

Windows 7, NVidia GeForce 425M.

我编写了一个简单的CUDA代码，该代码计算矩阵的行总和。矩阵具有一维表示形式（指向浮点数的指针）。

代码的串行版本如下（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）结果相等。

关键是，即使我更改了nThreadsPerBlock参数，CUDA版本几乎花费了计算串行时间两倍的时间：我测试了nThreadsPerBlock从32到1024（我的卡允许的每个块的最大线程数）。

IMO，矩阵尺寸大足以证明并行化：90,000 x 1,000。

下面，我报告使用different的串行和并行版本所花费的时间nThreadsPerBlock。报告msec的平均100样本时间为：

矩阵：nrow = 90000 x ncol = 1000

串行：每次采样的平均时间经过的毫秒（以100样品）289.18。

CUDA（32ThreadsPerBlock）：平均时间消逝每样毫秒（在100样本）497.11。

CUDA（1024ThreadsPerBlock）：平均时间消逝每样毫秒（在100样本）699.66。

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

我知道从主机复制到设备以及以其他方式进行复制时会产生某种开销，但是可能速度较慢是因为我没有实现最快的代码。

由于我远非CUDA专家：

我是否为此任务编写了最快的版本？如何改善我的代码？我可以摆脱内核函数中的循环吗？

任何想法表示赞赏。

编辑1

虽然我描述了一个标准rowSum，我有兴趣在AND/ OR具有行操作(0;1}的值，比如rowAND/ rowOR。就是说，正如一些评论员所建议的那样，它不允许我利用“ cuBLAS乘以1”的COL列向量技巧。

编辑2

根据用户的建议，其他用户在这里认可：

忘记尝试编写自己的功能，而是使用Thrust库，魔力来了。

Answer 1

如果这是您需要使用此数据进行操作的范围（汇总行），那么我预计GPU不会带来可观的收益。每个数据元素只有一个算术运算，为此您要付出将数据元素传输到GPU的费用。除了一定的问题大小（无论机器忙什么），由于算术强度为O（n），因此无法从更大的问题大小中获得更多好处。因此，这不是在GPU上解决的特别令人兴奋的问题。但是，正如前言所表明的那样，您在制作工艺上存在一个合并问题，这将进一步降低速度。让我们看一个小例子：&nbsp; &nbsp; C1&nbsp; C2&nbsp; C3&nbsp; C4R1&nbsp; 11&nbsp; 12&nbsp; 13&nbsp; 14R2&nbsp; 21&nbsp; 22&nbsp; 23&nbsp; 24R3&nbsp; 31&nbsp; 32&nbsp; 33&nbsp; 34R4&nbsp; 41&nbsp; 42&nbsp; 43&nbsp; 44上面是矩阵一小部分的简单图示示例。机器数据存储应将元素（11），（12），（13）和（14）存储在相邻的存储位置中。对于合并访问，我们需要一种访问模式，以便从同一条指令中请求相邻的内存位置，并在扭曲中执行。我们需要从warp的角度考虑代码的执行，即warp地执行32个线程。您的代码在做什么？在每个步骤/指令中都检索（要求）哪些元素？让我们看一下这行代码：&nbsp; &nbsp; &nbsp; &nbsp; sum+=m[rowIdx*ncol+k];rowIdx创建变量时，经线中的相邻线程具有相邻（即连续）的值。因此，当k= 0时，当我们尝试检索值时，每个线程都在请求哪个数据元素m[rowIdx*ncol+k]？在块0中，线程0的a rowIdx为0。线程1的a rowIdx为1，依此类推。因此，每个线程在此指令中要求的值是：Thread:&nbsp; &nbsp;Memory Location:&nbsp; &nbsp; Matrix Element:&nbsp; &nbsp; &nbsp;0&nbsp; &nbsp; &nbsp; m[0]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp;(11)&nbsp; &nbsp; &nbsp;1&nbsp; &nbsp; &nbsp; m[ncol]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; (21)&nbsp; &nbsp; &nbsp;2&nbsp; &nbsp; &nbsp; m[2*ncol]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; (31)&nbsp; &nbsp; &nbsp;3&nbsp; &nbsp; &nbsp; m[3*ncol]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; (41)但这不是合并访问！元素（11），（21）等在内存中不相邻。对于合并访问，我们希望“矩阵元素”行的内容如下：Thread:&nbsp; &nbsp;Memory Location:&nbsp; &nbsp; Matrix Element:&nbsp; &nbsp; &nbsp;0&nbsp; &nbsp; &nbsp; m[?]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp;(11)&nbsp; &nbsp; &nbsp;1&nbsp; &nbsp; &nbsp; m[?]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp;(12)&nbsp; &nbsp; &nbsp;2&nbsp; &nbsp; &nbsp; m[?]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp;(13)&nbsp; &nbsp; &nbsp;3&nbsp; &nbsp; &nbsp; m[?]&nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp; &nbsp;(14)如果您随后进行反向操作以确定?应该是什么值，那么您将得出如下指示：&nbsp; &nbsp; &nbsp; &nbsp; sum+=m[k*ncol+rowIdx];这将提供合并的访问权限，但不会为您提供正确的答案，因为我们现在正在汇总矩阵列而不是矩阵行。我们可以通过将您的数据存储重新组织为列优先顺序而不是行优先顺序来解决此问题。（您应该可以在Google上搜索到它的想法，对吗？）从概念上讲，这等效于转换矩阵m。如我所见，这是否方便您在我的问题范围之外，而实际上不是CUDA问题。在主机上创建矩阵或将矩阵从主机传输到设备时，这可能对您来说很简单。但总而言之，如果矩阵以行优先顺序存储，我不知道用100％合并访问来对矩阵行求和的方法。（您可以采用一系列的行减少操作，但这对我来说很痛苦。）当我们正在考虑在GPU上加速代码的方式时，考虑重新组织数据存储以方便GPU的情况并不少见。这是一个例子。而且，是的，我在这里概述的内容仍然在内核中保留了一个循环。作为补充说明，我建议分别对数据复制部分和内核（计算）部分进行计时。从您的问题中我无法确定您是在计时内核还是整个（GPU）操作，包括数据副本。如果单独对数据复制计时，则可能会发现仅数据复制时间超过了CPU时间。优化CUDA代码所做的任何努力都不会影响数据复制时间。在花费大量时间之前，这可能是有用的数据点。