CUDA中的非正方形矩陣乘法

我在CUDA中用於矩陣乘法的代碼可以讓我乘以正方形和非正方形矩陣，但寬度和高度必須是塊大小的倍數。因此，例如，我可以乘[3] [6] * [6] [3]（使用塊大小= 3），但我不能乘以[3] [2] * [2] [3 ]。CUDA中的非正方形矩陣乘法

有沒有人知道一種方法來做到這一點？這是我的核心：

#include <stdio.h> 

#include <limits.h> 

#include <stdlib.h> 
#define blocksize 3 
#define HM (1*blocksize) 
#define WM (2*blocksize) 
#define WN (1*blocksize) 
#define HN WM 
#define WP WN 
#define HP HM 
#define PTH WM 
#define PTW HM 

__global__ void nonsquare(float*M, float*N, float*P, int uWM,int uWN) 

{ 
__shared__ float MS[blocksize][blocksize]; 
__shared__ float NS[blocksize][blocksize]; 


int tx=threadIdx.x, ty=threadIdx.y, bx=blockIdx.x, by=blockIdx.y; 
int rowM=ty+by*blocksize; 
int colN=tx+bx*blocksize; 
float Pvalue=0; 


for(int m=0; m< uWM/blocksize;++m){ 
    MS[ty][tx]=M[rowM*uWM+(m*blocksize+tx)]; 
    NS[ty][tx]=M[colN + uWN*(m*blocksize+ty)]; 
    __syncthreads(); 

    for(int k=0;k<blocksize;k++) 
     Pvalue+=MS[ty][k]*NS[k][tx]; 
    __syncthreads(); 
    P[rowM*WP+colN]=Pvalue; 
    } 
    }

在此先感謝！

來源

2011-03-28 Bernardo

我認爲最簡單的辦法是隻墊上結尾的塊以零：

for(int m=0; m< uWM/blocksize;++m){ 
    colM = m*blocksize+tx; 
    rowN = m*blocksize+ty; 
    if (rowM > uWN || rowN > uWM || colM > uWM || colN > uWN) { 
     MS[ty][tx]=0.; 
     NS[ty][tx]=0.; 
    } else { 
     MS[ty][tx]=M[rowM*uWM+colM]; 
     NS[ty][tx]=N[colN + uWN*rowN]; 
    }

加減。（這NS行應引用N，不男，對吧？）

但是，因爲我似乎這裏是唯一一個使用現有的調庫時可能主張 - 爲什麼不使用CUBLAS或MAGMA而不是滾動您自己？它們速度很快，並經過數百名用戶的測試。

來源

2011-03-28 15:28:07

這似乎是一個好點，雖然因爲我使用矩陣代表醫學圖像，我不確定如果用零填充將影響圖像。但我會試一試。至於使用CUBLAS，我一直在閱讀它，但我有點擔心它有很多新的語法（學習曲線高嗎？）。但是我會盡一切努力的。感謝您的回覆;） – Bernardo 2011-03-28 16:16:39

陣列只在矩陣乘法例程中被填充。它們不會被傳回，並且它們不會影響最終結果，因爲您只是向矩陣元素添加零。至於CUBLAS（或岩漿，或其他） - 學習曲線是真實的，但事後你不必編寫自己的線性代數例程，而且性能是一個賣點...... – 2011-03-28 17:29:43

我明白了......我將很快處理的事情之一是非常大（不適合GPU卡存儲器）3D矩陣，CUBLAS庫可以消除3D索引和內存分配的痛苦（我認爲第一個問題是比第二個更可行）？ – Bernardo 2011-03-28 17:41:20

這裏的基本性能要求是，共享內存「tile」的第一維或第二維必須是16的整數倍 - 歷史上，這是實現最佳全局內存帶寬所必需的（對於計算能力爲1.3半經合併交易）。是否應該是瓦片的第一個或第二個維度取決於矩陣是以列還是行的主要順序存儲。沒有什麼可說共享內存瓦需要是方形的，只是存儲的主要維度（BLAS表示法中的LDA）是16的整數倍。

您可以很容易地將內核的瓦片尺寸模板參數並根據矩陣維度實例化幾個版本。對於給定的體系結構，可能存在一個平衡佔用率和指令級並行度的最佳平鋪維度。解決這個問題的「巧妙」方法很可能是將矩陣乘法分解成兩個操作 - 第一個以最優分塊大小完成大部分工作，第二個分配大小不同的剩餘列。如果結果在產品完成後直接返回到主機內存，則最好在主機上使用與GPU內核重疊的優化的BLAS完成第二個操作。這是UTK Magma庫中許多例程使用的方法。

來源

2011-03-29 10:29:21 talonmies

CUDA中的非正方形矩陣乘法

回答

相關問題