大型矩阵的 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() 来仔细检查，问题似乎出在 cudaThreadSynchronize() 而不是 cudaMemcpy ()。

更改 TDR 超时后，您是否重新启动了机器？遗憾的是，需要重新启动 Windows 才能更改 TDR 设置。 This Microsoft document对可用的完整设置有相当好的描述。

内核问题

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

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

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

但是当你启动内核时，你没有指定为每个 block 分配多少共享内存:

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

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

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

主要问题

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

我应该早点建议，但是运行 cuda-memcheck 来检查这样的内存访问冲突通常很有用。