CUDAマトリックスの乗算は、大きなマトリックスのために破損します

Question

CUDA 3.2とVS 2008を使用して実装された次のマトリックス乗算コードがあります。WindowsServer2008 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<<<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タイムアウトを変更した後、マシンを再起動しましたか？残念ながら、TDR設定を変更するには、Windowsを再起動する必要があります。 このMicrosoftドキュメント 利用可能な完全な設定についてかなり良い説明があります。

カーネルの問題

この場合、問題は実際にはWDDMタイムアウトではありません。カーネルには解決する必要があるエラーがあります（たとえば、増分できるはずです i 反復ごとに複数）とチェックアウト matrixMul SDKのサンプルが役立つ場合があります。ちなみに、実際には（パフォーマンスのために）Cublasを使用してマトリックスの乗算を実行する方が良いので、これが学習演習であることを願っています。

コードの最も重要な問題は、実際に割り当てることなく共有メモリを使用していることです。あなたのカーネルにあなたは持っています：

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

ただし、カーネルを起動すると、各ブロックに割り当てる共有メモリの量を指定しません。

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

<<< >>>構文は、実際には3番目と4番目がオプションである4つの引数を取ります。 4番目は、計算とデータ転送の間の重複を取得するために使用されるストリームインデックスです（および同時カーネルの実行）が 第3 引数は、ブロックごとに共有メモリの量を指定します。この場合、私はあなたが保存したいと思います TileWidth * TileWidth 共有メモリにフロートするので、使用します。

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

主な問題

コメントで言及しているように、実際の問題は、マトリックスの幅がブロック幅の倍数ではなかったことです（および高さは正方形であるためです。つまり、端を超えたスレッドは配列の端を超えてアクセスします。コードはどちらかです。非倍数のケースを処理するか、幅がブロックサイズの倍数であることを確認する必要があります。

私はこれを早く提案すべきだったが、それはしばしば走るのに役立つ cuda-memcheck このようなメモリーアクセス違反を確認する。

Answer 2

ドライバーのタイムアウト設定を変更する必要があります。これは、システムを反応させないドライバーの故障を防ぐためのWindows機能です。を確認します マイクロソフトページ それを行う方法を説明します。

Answer 3

また、GPUデバイスの「タイムアウト」フラグ設定を確認する必要があります。 Cuda SDKをインストールしている場合、「DeviceQuery」アプリがこのプロパティを報告すると思います。