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已经返回了一个错误？ （I证实它没有在调试时）。

任何想法/帮助将不胜感激！......非常感谢球员！

内核代码：

//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设置。 这个微软文档拥有的全部可用设置一个相当不错的描述。

<强> 内核问题

在这种情况下，问题是不实际的WDDM超时。有些情况下，你将需要解决（比如你应该能够超过一个在每次迭代incremement 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功能，防止故障的驱动程序，使系统反应迟钝。 检查微软页面描述如何做到这一点。

Answer 3

您也应该检查你的GPU设备的“超时”标志设置。如果你安装了CUDA SDK，我相信“DEVICEQUERY”应用程序会报告这个属性。