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

Posted 2023-04-15

【中文标题】大型矩阵的 CUDA 矩阵乘法中断

【英文标题】：CUDA Matrix multiplication breaks for large matrices

【发布时间】：2011-05-02 20:50:03

【问题描述】：

我有以下矩阵乘法代码，使用 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;

【问题讨论】：

能发一下内核代码吗？

编辑：添加了两个内核代码函数

为了完整起见，您能否发布 dimGrid/dimBlock 分配？我假设dimGrid.x = dimGrid.y 和dimBlock.x = dimBlock.y，而那个Width 是dimBlock.x 的倍数？

@Tom，我添加了我的 dimGrid 和 dimBlock 定义并在那里发现了问题.. 我没有使用多个 dimBlock... 我的内核代码没有正确处理.. 谢谢你很多!!!!，请将此评论添加到下面的答案中，以便我选择它。再次感谢！！！

不客气。我应该建议早点运行cuda-memcheck，对此感到抱歉。它应该立即检测到越界访问。

【参考方案1】：

控制 WDDM 超时

问题实际上是内核而不是cudaMemcpy()。当您启动内核时，GPU 会关闭并与 CPU 异步执行工作，因此只有当您与 GPU 同步时，您才必须等待工作完成。 cudaMemcpy() 涉及隐式同步，因此您会看到问题所在。

您可以通过在内核之后调用cudaThreadSynchronize() 来仔细检查这一点，问题似乎出在cudaThreadSynchronize() 而不是cudaMemcpy()。

更改 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);

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

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

主要问题

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

我应该早些提出这个建议，但运行cuda-memcheck 来检查内存访问违规通常很有用。

【讨论】：

我确实重新启动并且更改生效，因为黑屏出现的时间更长..但是它仍然崩溃....

我回家后会尝试 cudaThreadSynchronize 并重新发布！

好吧，你是对的。我刚刚这样做了，我得到了同样的错误......我将 TdrDelay 作为 REG_DWORD 添加到 HKEY_LOCAL_MACHINE\SYSTEM\CurrentControlSet\Contol\GraphicsDrivers。我重新启动了我的机器，我注意到屏幕变黑并出现错误需要更长的时间..只要我设置了延迟..但它仍然不起作用。我不完全相信这是一个延迟，因为它处理 2500 的宽度就好了，但除此之外，它崩溃了.. 甚至 2800 ......我错过了什么吗？

我现在还注意到，当我添加 cudaThreadSynchronize() 时，它会破坏它以前用于工作的大小矩阵... 2500 会在添加该行时破坏它。 （我不知道这是否给你们任何线索）

IIUC 高达 ~2500 它可以快速完成（比如不到几秒钟），但超过此时间，无论您将 TDR 设置多高，它都不会完成。您可以通过将 TDR 设置得很高来检查这一点，请耐心等待！既然如此，下一步就是看内核了。

【参考方案2】：

您必须更改驱动程序超时设置，是 Windows 功能，以防止有故障的驱动程序使系统无响应。 查看描述如何执行此操作的 Microsoft Page 。

【讨论】：

我应该尝试除 TdrDelay 之外的其他方法吗？

【参考方案3】：

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

【讨论】：

感谢您的回复！我该去哪里修改这个属性？

我不确定如何修改它 - 这是驱动程序处理的事情。这可能与您是否将显示器连接到设备有关。