对于大型矩阵，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”应用会报告这个属性。