使用CUDA进行动态矩阵乘法

Question

使用CUDA进行动态矩阵乘法

Dan*_*Dan 2 cuda matrix matrix-multiplication

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

我希望将输入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;


}

Run Code Online (Sandbox Code Playgroud)

更新:我决定采用零填充.但是我的答案不正确.将矩阵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

Run Code Online (Sandbox Code Playgroud)

矩阵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

Run Code Online (Sandbox Code Playgroud)

所以CI获得的结果是正确的:

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

Run Code Online (Sandbox Code Playgroud)

但是,如果你去除零,矩阵应该是:A:

5.000 0.000
9.000 0.000

Run Code Online (Sandbox Code Playgroud)

B:

7.000 4.000
8.000 7.000

Run Code Online (Sandbox Code Playgroud)

C应该是:

35.000 20.000
63.000 36.000

Run Code Online (Sandbox Code Playgroud)

然而,两个矩阵Cs不相同.

Answer 1

tal*_*ies 7

这不是一个非常明确的问题,所以这个答案可以根据你之前在几个相似的问题中提出的问题进行猜测.

理解如何进行这种操作的一个很好的起点是回到起点并从第一原理考虑矩阵 - 矩阵乘法问题.您感兴趣的代码是计算两个矩阵的点积,C = AB.您所拥有的限制是您使用的内核只能计算矩阵的乘积,这些矩阵是某个内部块大小的倍数.所以,你可以做什么？

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

在此输入图像描述

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

在此输入图像描述

可能不会立即明白这有助于解决问题,但让我们考虑一个具体的例子:

您有一个最佳矩阵乘法内核,它使用32的内部块大小,并且仅当矩阵是该块大小的圆倍数时才是正确的.
你有一对1000x1000平方矩阵可以乘以.

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

如果您决定使用1024x1024产品,则可以使用块分解思想来制定如下问题:

在此输入图像描述

其中O _nn表示适当大小的零矩阵.现在你有了一对1024x1024矩阵,它们的产品就会产生

在此输入图像描述

即.左手,上部块是包含AB的1000x1000矩阵.这实际上是零填充以实现正确的结果.在这个例子中,它意味着执行的计算比所需的多7%.是否重要可能是应用程序特定的.

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

在此输入图像描述

与甲₁₁个乙₁₁是992x992矩阵,ö _NN是零点矩阵如前.在第一次检查时,这看起来并不是很有帮助,但值得记住的是,使右侧矩阵的所有计算仅包含计算矩阵乘积所需的总计算量的约1.2%.在GPU进行主要计算时,可以在主机CPU上轻松完成它们,然后将其添加到GPU结果中以形成最终矩阵.由于CUDA API是异步的,因此大多数主机计算可以完全隐藏并且实际上是免费的.

这个答案包含两个策略,用于执行您要求的内容,而不会更改当前内核代码的单行.显然有第三种方法,即更彻底地修改内核本身,但这是你应该首先尝试自己,然后在你的解决方案不起作用时寻求帮助的方法.

归档时间：	13 年，7 月前
查看次数：	1766 次
最近记录：	13 年，7 月前