CUDA矩阵乘法中断了大型矩阵

我有以下矩阵乘法代码,使用CUDA 3.2和VS 2008实现。我在Windows server 2008 r2 enterprise上运行。 我正在运行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<<>>(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<<>>(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; } 

控制WDDM超时

问题实际上是内核而不是cudaMemcpy() 。 当您启动内核时,GPU会关闭并与CPU异步完成工作,因此只有在与GPU同步时才需要等待工作完成。 cudaMemcpy()涉及隐式同步,因此您可以在此处查看问题。

您可以通过在内核之后调用cudaThreadSynchronize()仔细检查这个问题,问题将显示在cudaThreadSynchronize()而不是cudaMemcpy()

更改TDR超时后,您是否重新启动了计算机? 不幸的是,Windows需要重新启动才能更改TDR设置。 此Microsoft文档对可用的完整设置有相当好的描述。

内核问题

在这种情况下,问题实际上不是WDDM超时。 内核中存在您需要解决的错误(例如,您应该能够在每次迭代时将i递增多于一个)并且检查SDK中的matrixMul示例可能很有用。 顺便说一下,我希望这是一个学习练习,因为实际上使用CUBLAS执行矩阵乘法会更好(性能)。

代码中最关键的问题是您使用共享内存而不实际分配任何内存。 在你的内核中你有:

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

但是当您启动内核时,您没有指定为每个块分配多少共享内存:

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

<<< >>>语法实际上需要四个参数,其中第三个和第四个是可选的。 第四个是流索引,用于获取计算和数据传输之间的重叠(以及并发内核执行),但第三个参数指定每个块的共享内存量。 在这种情况下,我假设您要在共享内存中存储TileWidth * TileWidth浮点数,因此您将使用:

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

主要问题

正如您在评论中提到的,实际问题是您的矩阵宽度不是块宽度的倍数(和高度,因为它是正方形,这意味着超出末尾的线程将访问超出数组末尾。代码应该是处理非多重情况或应确保宽度是块大小的倍数。

我应该早些时候提出这个建议,但运行cuda-memcheck来检查像这样cuda-memcheck访问违规通常很有用。

您必须更改驱动程序超时设置,Windowsfunction是为了防止错误的驱动程序使系统无响应。 检查描述如何执行此操作的Microsoft页面 。

您还应该检查GPU设备上的“超时”标志设置。 如果您安装了CUDA SDK,我相信“deviceQuery”应用程序将报告此属性。