matrix - 使用 CUDA 进行动态矩阵乘法

标签 matrix matrix-multiplication cuda

我一直试图编写的简单程序的想法是从用户那里获取输入以查看要乘以多大的矩阵。

我希望将输入 x 乘以 x，我目前不希望将两种不同的大小相乘。

你们建议我如何完成这项工作？

对不起，我的问题不够清楚，我想修改这个内核，以便它可以处理任何大小的矩阵(其中 x 和 y 是等价的，以保持简单)。而不是 16 的倍数。

我不确定您是否需要我当前的代码，但这是内核代码:

// CUDA Kernel
__global__ void matrixMul( float* C, float* A, float* B, int wA, int wB,size_t block_size)
{
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int tx = threadIdx.x;
    int ty = threadIdx.y;

    int aBegin = wA * block_size * by;
    int aEnd   = aBegin + wA - 1;
    int aStep  = block_size;

    int bBegin = block_size * bx;

    int bStep  = block_size * wB;
    float Csub=0;

    for (int a = aBegin, b = bBegin; a <= aEnd; a += aStep, b += bStep) 
    {
        extern __shared__ float As[];
        extern __shared__ float Bs[];
        extern __shared__ float smem[];

        smem[ty*block_size+tx] = A[a + wA * ty + tx];

        smem[block_size*block_size+ty*block_size+tx]  = B[b + wB * ty + tx];

        __syncthreads();

        for (int k = 0; k < block_size; ++k)
            Csub += smem[ty*block_size+k] * smem[block_size*block_size+k*block_size+tx] ;

        __syncthreads();
    }

    int c = wB * block_size * by + block_size * bx;
    C[c + wB * ty + tx] = Csub;


}

更新:我决定使用零填充。但是我得到不正确的答案。 取矩阵 A 2x2，填充到 16x16:

5.000 0.000 9.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000

矩阵 B，2x2 填充到 16x16:

7.000 4.000 8.000 7.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000
0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000 0.000

所以我得到的 C 的结果是正确的:

35.000 20.000 40.000 35.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000
 0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000  0.000

然而，如果你去掉零，矩阵应该是: 答:

5.000 0.000
9.000 0.000

乙:

7.000 4.000
8.000 7.000

C 应该是:

35.000 20.000
63.000 36.000

然而，这两个矩阵 C 并不相同。

最佳答案

这不是一个非常明确的问题，因此这个答案是基于您之前在几个相当相似的问题中提出的问题的猜测。

要理解如何进行此类运算，一个很好的起点是回到起点，从第一性原理考虑矩阵-矩阵乘法问题。您对计算两个矩阵 C = AB 的点积的代码感兴趣。您的限制是您使用的内核只能计算矩阵的乘积，这些矩阵是某个内部 block 大小的整数倍。那你能做什么？

看待这个问题的一种方法是想象A 和B 矩阵是block matrices .矩阵乘法可以这样写:

enter image description here

然后可以通过 A 和 B 中的八个子矩阵的乘积组合形成生成的矩阵 C:

enter image description here

这对解决问题的帮助可能不是很明显，但让我们考虑一个具体的例子:

  1. 您有一个最佳矩阵乘法内核，它使用内部 block 大小 32，并且仅当矩阵是该 block 大小的整数倍时才正确。
  2. 你有一对 1000x1000 的方阵要相乘。

这些第一个事实意味着您的内核只能正确解决 1024x1024 产品或 992x992 产品，但不能解决您需要的 1000x1000 操作。

如果你决定使用 1024x1024 的产品，你可以使用 block 分解的思想来制定这样的问题:

enter image description here

其中 Onn 表示适当大小的零矩阵。现在您有一对 1024x1024 矩阵，它们的乘积将导致

enter image description here

即。左手边，上面的 block 是一个包含 AB 的 1000x1000 矩阵。这实际上是零填充以实现正确的结果。在此示例中，这意味着执行的计算比所需的多大约 7%。重要与否可能取决于具体应用。

第二种方法是使用基本内核计算一个 992x992 产品，然后制定一个策略来处理 block 分解版本计算中的其他七个产品，如下所示:

enter image description here

A11 和 B11 是 992x992 矩阵，O nn 和以前一样是零矩阵。乍一看，这看起来不是很有帮助，但值得记住的是，所有 使右侧矩阵的计算仅包含计算矩阵乘积所需的总计算的大约 1.2%。它们可以在 GPU 进行主要计算时在主机 CPU 上轻松完成，然后添加到 GPU 结果以形成最终矩阵。因为 CUDA API 是异步的，所以大部分主机计算都可以完全隐藏并且实际上是免费的。

此答案包含两种策略，可以只更改当前内核代码的一行 来完成您的要求。显然还有第三种方法，即更彻底地修改内核本身，但这是您应该首先自己尝试的方法，如果您的解决方案不起作用，则寻求帮助。

关于matrix - 使用 CUDA 进行动态矩阵乘法，我们在Stack Overflow上找到一个类似的问题： https://stackoverflow.com/questions/9250897/

上一篇：ruby-on-rails - Rails 生产服务器缓存权限被 Nginx 拒绝

下一篇：sql - 简单的sql : how do I group into separate columns?

相关文章：

python - 为 numpy 数组的每个 "other"字段 ((i+j)%2==0) 添加值

c++ - CUDA 在使用函数指针时将主机函数作为内核启动

matlab - 如何以异步方式有效地将变量从 Matlab 传递到 GPU？

matlab - Matlab/CUDA : ocean wave simulation

matlab - 将 n 个长度为 p 的向量乘以 n 个大小为 pxp 的矩阵

python - 矩阵乘法问题 - Numpy vs Matlab？

python - tensorflow 中的张量形状

python - 值错误 : matrices are not aligned

documentation - 是否有关于 Blitz++ 矩阵的文档？

c - C 中的矩阵乘法无法正常工作

©2024 IT工具网  联系我们