与CUDA的动态矩阵乘法

我一直试图写的简单程序的想法是从用户那里获取输入，以查看矩阵的多大乘法。与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; 


}

更新：我决定去填零。但是我得到不正确的答案。以矩阵2×2，填充为16×16：

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，2×2填充为16×16：

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

所以对于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

但是，如果你去掉矩阵应该是的零点： - 答：

5.000 0.000 
9.000 0.000

B：

7.000 4.000 
8.000 7.000

C应该：

35.000 20.000 
63.000 36.000

然而两个矩阵Cs的是不一样的。

来源

2012-02-12 Dan

你的问题是什么？你在问如何从用户那里获得输入？ – harrism 2012-02-13 04:09:02

如果我正确理解了你以前的问题，你真正的问题是如何修改这个内核代码（它本身就是CUDA SDK矩阵乘法例子的一个很轻微的修改版），所以它可以用来乘以任意大小的矩阵，而不是圆内核块大小的倍数。你能编辑你的问题来反映这个吗？目前，你真的不知道你在问什么。 – talonmies 2012-02-13 10:37:15

@talonmies，你是对的。那正是我在寻找的 – Dan 2012-02-13 14:47:33

这不是一个非常明确的问题，所以这个答案是基于你以前在几个相似的问题中提出的问题。

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

一种方式来看待这个问题的方法是想象一个和乙矩阵是block matrices。的矩阵乘法可以被写成这样：

enter image description here

和然后可以通过由八个子矩阵的产品的组合中形成和乙所得矩阵C：

它可能不是马上就会明白这是如何帮助解决问题的，但让我们考虑一个具体的例子：

您有一个使用32的内部块的大小，并且只有当矩阵是块大小的倍数轮正确的最佳矩阵乘法内核。

你有一对1000×1000方阵来乘。

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

如果你决定使用1024x1024的产品，你可以使用块分解理念，制定这样的问题：

其中Ø_NN表示零的合适尺寸的矩阵。现在你有一双1024x1024矩阵，以及他们的产品会导致

即。左上方块是包含AB的1000x1000矩阵。这实际上是零填充以实现正确的结果。在这个例子中，这意味着执行的计算量比所需的多7％。无论重要与否都可能是特定于应用程序的。

第二种方法是使用的基本内核计算一个992x992的产品，然后制定出一项战略，以应对其他七种产品在计算块分解的版本，像这样：

与甲和乙为992x992矩阵，ö_ン与以前一样是零矩阵。在第一次检查时，这看起来不是很有用，但值得记住的是，使右侧矩阵仅包含计算矩阵乘积所需总计算量的约1.2％的计算。在GPU进行主计算时，他们可以在主CPU上轻松完成，然后添加到GPU结果中以形成最终矩阵。由于CUDA API是异步的，因此大部分主机计算都可以完全隐藏并且实际上是免费的。

这个答案包含两个策略，用于做你要求的，而不会改变你当前内核代码的单行内容。显然还有第三种方式，这是更彻底地修改内核本身，但这是你应该先尝试自己，然后寻求帮助，如果你的解决方案不起作用。