對於大型矩陣，CUDA矩陣乘法中斷

Question

我使用CUDA 3.2和VS 2008實現了以下矩陣乘法代碼。我在Windows Server 2008 R2企業版上運行。我正在運行Nvidia GTX 480.以下代碼可以很好地處理「寬度」（矩陣寬度）值高達2500左右的值。

int size = Width*Width*sizeof(float); 
float* Md, *Nd, *Pd; 
cudaError_t err = cudaSuccess; 

//Allocate Device Memory for M, N and P 
err = cudaMalloc((void**)&Md, size); 
err = cudaMalloc((void**)&Nd, size); 
err = cudaMalloc((void**)&Pd, size); 

//Copy Matrix from Host Memory to Device Memory 
err = cudaMemcpy(Md, M, size, cudaMemcpyHostToDevice); 
err = cudaMemcpy(Nd, N, size, cudaMemcpyHostToDevice); 

//Setup the execution configuration 
dim3 dimBlock(TileWidth, TileWidth, 1); 
dim3 dimGrid(ceil((float)(Width)/TileWidth), ceil((float)(Width)/TileWidth), 1); 

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width); 

err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); 

//Free Device Memory 
cudaFree(Md); 
cudaFree(Nd); 
cudaFree(Pd); 

當我設置的「寬」到3000或更高，我得到一個黑色的屏幕後，出現以下錯誤：

我在網上看了一下，我看到有些人有這個問題，因爲看門狗在掛起超過5秒後殺死內核。我試圖在註冊表中編輯「TdrDelay」，這會延遲黑屏和同樣錯誤出現之前的時間。所以我認爲這不是我的問題。

我調試到我的代碼，發現此行是罪魁禍首：

err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost); 

這是我用回我的結果從設備設置我的矩陣乘法內核函數被調用後。一直到這一點似乎運行良好。我相信我正確地分配內存，並不知道爲什麼會發生這種情況。我想也許我沒有足夠的內存在我的卡上，但不應該cudaMalloc返回錯誤？ （我確認它沒有在調試時）。

任何想法/援助將不勝感激！...謝謝很多傢伙！

內核代碼：

//Matrix Multiplication Kernel - Multi-Block Implementation 
__global__ void MatrixMultiplicationMultiBlock_Kernel (float* Md, float* Nd, float* Pd, int Width) 
{ 
int TileWidth = blockDim.x; 

//Get row and column from block and thread ids 
int Row = (TileWidth*blockIdx.y) + threadIdx.y; 
int Column = (TileWidth*blockIdx.x) + threadIdx.x; 

//Pvalue store the Pd element that is computed by the thread 
float Pvalue = 0; 

for (int i = 0; i < Width; ++i) 
{ 
    float Mdelement = Md[Row * Width + i]; 
    float Ndelement = Nd[i * Width + Column]; 
    Pvalue += Mdelement * Ndelement; 
} 

//Write the matrix to device memory each thread writes one element 
Pd[Row * Width + Column] = Pvalue; 
} 

我也有使用共享內存此等功能，同時也給出了同樣的錯誤：

電話：

  MatrixMultiplicationSharedMemory_Kernel<<<dimGrid, dimBlock, sizeof(float)*TileWidth*TileWidth*2>>>(Md, Nd, Pd, Width); 

內核代碼：

//Matrix Multiplication Kernel - Shared Memory Implementation 
__global__ void MatrixMultiplicationSharedMemory_Kernel (float* Md, float* Nd, float* Pd, int Width) 
{ 
int TileWidth = blockDim.x; 

//Initialize shared memory 
extern __shared__ float sharedArrays[]; 
float* Mds = (float*) &sharedArrays; 
float* Nds = (float*) &Mds[TileWidth*TileWidth]; 

int tx = threadIdx.x; 
int ty = threadIdx.y; 

//Get row and column from block and thread ids 
int Row = (TileWidth*blockIdx.y) + ty; 
int Column = (TileWidth*blockIdx.x) + tx; 
float Pvalue = 0; 

//For each tile, load the element into shared memory 
for(int i = 0; i < ceil((float)Width/TileWidth); ++i) 
{ 
    Mds[ty*TileWidth+tx] = Md[Row*Width + (i*TileWidth + tx)]; 
    Nds[ty*TileWidth+tx] = Nd[(ty + (i * TileWidth))*Width + Column]; 

    __syncthreads(); 

    for(int j = 0; j < TileWidth; ++j) 
    { 
     Pvalue += Mds[ty*TileWidth+j] * Nds[j*TileWidth+tx]; 
    } 

    __syncthreads(); 
} 

//Write the matrix to device memory each thread writes one element 
Pd[Row * Width + Column] = Pvalue; 
} 

Answer 1

控制WDDM超時

這個問題實際上是內核不是cudaMemcpy()。當您啓動內核時，GPU會關閉並與CPU異步執行工作，所以只有在與GPU同步時才需要等待工作完成。 cudaMemcpy()涉及隱式同步，因此這就是您看到問題的地方。

在內核和問題出現在cudaThreadSynchronize()而不是cudaMemcpy()之後，您可以通過調用cudaThreadSynchronize()來仔細檢查此問題。

更改TDR超時後，您是否重新啓動機器？不幸的是Windows需要重新啓動才能更改TDR設置。 This Microsoft document對可用的完整設置有相當好的描述。

核心問題

在這種情況下，問題是不實際的WDDM超時。內核中有需要解決的錯誤（例如，您應該可以在每次迭代中將i增加一個以上），並檢查SDK中的matrixMul示例可能會有用。順便提一句，我希望這是一個學習練習，因爲實際上使用CUBLAS執行矩陣乘法會更好（表現）。

代碼中最關鍵的問題是您正在使用共享內存而不實際分配任何內存。在你的內核，你有：

//Initialize shared memory 
extern __shared__ float sharedArrays[]; 

但是當你進入內核不指定多少共享內存來分配每個塊：

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width); 

的< < < >>>語法，其實需要四個參數，其中第三個和第四個是可選的。第四個是用於在計算和數據傳輸（以及併發內核執行）之間獲得重疊的流索引，但是參數指定每個塊的共享內存數量。在這種情況下，我認爲你要存儲TileWidth * TileWidth漂浮在共享內存，因此會使用：

MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock, dimBlock.x * dimBlock.x * sizeof(float)>>>(Md, Nd, Pd, Width); 

的主要問題

正如您在您的評論提到，實際問題是你的矩陣寬度不是塊寬度的倍數（和高度，因爲它是方形的，這意味着超出末端的線程將超出數組的末尾。代碼應該處理非多重情況，否則它應該確保寬度是塊大小的倍數。

我應該早些提出這個建議，但通常運行cuda-memcheck來檢查這樣的memeory訪問衝突是非常有用的。

Answer 2

你必須更改驅動程序超時設置，是Windows功能，以防止有故障的驅動程序，使系統無響應。 檢查Microsoft Page描述如何做到這一點。

Answer 3

您還應該檢查GPU設備上的「超時」標誌設置。如果你安裝了CUDA SDK，我相信「deviceQuery」應用會報告這個屬性。