CUDA矩阵乘法在大矩阵下出现故障

6

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

当我将“Width”设置为3000或更大时,黑屏后出现以下错误: screenshot 我在网上查找后发现,有些人之所以会遇到这个问题,是因为看门狗在挂起超过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;
}

请您发布核心代码好吗? - Tom
EDIT:添加了两个内核代码函数 - ntsue
1
请问您能否提供dimGrid/dimBlock的赋值代码?我假设dimGrid.x = dimGrid.y和dimBlock.x = dimBlock.y,并且Width是dimBlock.x的倍数,以保证完整性。 - Tom
@Tom,我添加了我的dimGrid和dimBlock定义,并在其中找到了问题...我没有使用dimBlock的倍数...我的内核代码没有正确处理这个问题...非常感谢你!!!请将这个评论添加到你下面的答案中,这样我就可以选择它了。再次感谢!!! - ntsue
不用谢。我应该早点建议运行cuda-memcheck,对此感到抱歉。它应该能立即检测到越界访问。 - Tom
3个回答

10

控制WDDM超时

问题实际上是内核而不是cudaMemcpy()。当您启动内核时,GPU会异步地与CPU一起工作,因此只有在与GPU同步时,您才必须等待工作完成。 cudaMemcpy() 包含隐式同步,因此这就是您看到问题的地方。

您可以通过在内核后调用cudaThreadSynchronize()来进行双重检查,这时问题将出现在 cudaThreadSynchronize() 而不是 cudaMemcpy() 上。

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

内核问题

在这种情况下,实际上问题并不是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检查这样的内存访问违规,这通常很有用。


好的,你是对的。我刚刚这样做了,我得到了相同的错误... 我将TdrDelay添加为REG_DWORD到HKEY_LOCAL_MACHINE\SYSTEM\CurrentControlSet\Contol\GraphicsDrivers。我重新启动了我的机器,我注意到屏幕变黑并出现错误需要更长的时间.. 大约和我设置的延迟一样长.. 但它仍然不起作用。我并不完全相信这是一个延迟问题,因为它可以很好地处理2500的宽度,但是任何比那更多的东西都会崩溃.. 即使是2800... 我错过了什么吗? - ntsue
如果我理解正确,当处理的数据量在2500左右时,它会很快完成(大约不到几秒钟),但是超过这个数量时,无论你将TDR设置得多高,它都永远无法完成。您可以通过将TDR设置得非常高来检查此问题,只需耐心等待即可!既然如此,下一步就是查看内核。 - Tom
谢谢Tom,实际上我发布了错误的签名..我有两个矩阵函数..一个优化共享内存,另一个不是..事实上,两者都无法处理大值..我将编辑以更新正确的函数调用..但实际上我确实分配了内存。 - ntsue
抱歉让你感到困惑 =/...我更新了帖子,实际上我分配了TileWidth * TileWidth * sizeof(float) * 2的空间...因为我有两个矩阵Nds和Mds用于存储我的值..这样应该没问题,对吧? - ntsue
是的,但现在让我们集中精力于非共享内存版本,因为它以相同的方式失败。最好不要过于复杂化! - Tom
显示剩余3条评论

1
你需要更改驱动程序超时设置,这是 Windows 功能,用于防止有问题的驱动程序导致系统无响应。 请查看 Microsoft 页面,了解如何操作。

除了TdrDelay,我应该尝试其他什么吗? - ntsue

0
你还应该检查GPU设备上的“timeout”标志设置。如果你已经安装了CUDA SDK,我相信“deviceQuery”应用程序会报告这个属性。

谢谢您的回复!我在哪里可以修改这个属性? - ntsue
我不确定如何修改它 - 这是驱动程序处理的事情。这可能与您是否连接了显示设备有关。 - Edric

网页内容由stack overflow 提供, 点击上面的
可以查看英文原文,
原文链接