與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是異步的，因此大部分主機計算都可以完全隱藏並且實際上是免費的。

這個答案包含兩個策略，用於做你要求的，而不會改變你當前內核代碼的單行內容。顯然還有第三種方式，這是更徹底地修改內核本身，但這是你應該先嚐試自己，然後尋求幫助，如果你的解決方案不起作用。