CUDA Matrixmultiplikation Pausen für große Matrizen

Question

Ich habe die folgende Matrixmultiplikation Code, implementiert mit CUDA 3.2 und VS 2008. Ich auf Windows Server 2008 R2 Enterprise leite. Ich arbeite mit einer Nvidia GTX 480. Der folgende Code funktioniert gut mit Werten von „Breite“ (Matrix Breite) bis etwa 2500 oder so.

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

Wenn ich die „Breite“ auf 3000 oder höher gesetzt, bekomme ich folgende Fehlermeldung nach einem schwarzen Bildschirm:

Ich sah online und ich sah, dass einige Leute dieses Problem hat, weil der Watchdog den Kernel tötete, nachdem er für mehr als 5 Sekunden hängt. Ich versuchte, die „TdrDelay“ in der Registrierung bearbeiten und dies verzögert die Zeit, bevor der schwarze Bildschirm und gleiche Fehler aufgetreten. Also schloss ich das nicht mein Problem war.

Ich gedebuggt in meinen Code und fand diese Linie die Täter zu sein:

err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);

Dies ist, was ich verwende, um meine Ergebnismenge aus dem Gerät zurück, nachdem meine Matrixmultiplikation Kernel-Funktion aufgerufen wird. Alles bis zu diesem Punkt gut zu laufen scheint. Ich glaube, ich bin die Zuweisung korrekt Speicher und kann nicht herausfinden, warum dies geschieht. Ich dachte, vielleicht habe ich nicht genug Speicher auf meiner Karte für diese haben, aber dann sollte nicht cudaMalloc haben einen Fehler zurückgegeben? (I bestätigte es nicht tat während des Debuggens).

Irgendwelche Ideen / Hilfe wäre sehr dankbar! ... Vielen Dank Jungs !!

Kernel-Code:

//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;
}

Ich habe auch diese andere Funktion, die verwendet den gemeinsamen Speicher, und es gibt auch den gleichen Fehler:

Call:

            MatrixMultiplicationSharedMemory_Kernel<<<dimGrid, dimBlock, sizeof(float)*TileWidth*TileWidth*2>>>(Md, Nd, Pd, Width);

Kernel-Code:

 //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

Die Steuerung der WDDM Timeout

Das Problem ist eigentlich der Kernel nicht die cudaMemcpy(). Wenn Sie den Kernel starten geht der GPU ab und macht die Arbeit asynchron mit der CPU, so dass es nur ist, wenn Sie mit der GPU zu synchronisieren, dass Sie für die Arbeit zu Ende warten. cudaMemcpy() beinhaltet eine implizite Synchronisation, also das ist, wo Sie das Problem zu sehen.

Sie könnten doppelt, um diese durch cudaThreadSynchronize() nach dem Kernel aufrufen und das Problem wird auf dem cudaThreadSynchronize() sein anstelle des cudaMemcpy().

Nach dem TDR-Timeout zu ändern, haben neu starten Sie Ihre Maschine? Leider muss Windows neu gestartet werden, um die TDR-Einstellungen zu ändern. Dieses Microsoft-Dokument eine ziemlich gute Beschreibung der vollständigen Einstellungen zur Verfügung hat.

Kernel Probleme

In diesem Fall ist das Problem nicht eigentlich das WDDM-Timeout. Es gibt Fehler im Kernel, die Sie zu lösen benötigen würden (zum Beispiel sollten Sie bei jeder Iteration um mehr als ein bis incremement i der Lage sein) und die matrixMul Probe in der SDK-Check-out nützlich sein können. Übrigens, ich hoffe das eine Lernübung ist, da in der Realität wäre es besser, (für die Leistung) CUBLAS mit Matrixmultiplikation auszuführen.

Das kritischste Problem in dem Code ist, dass Sie gemeinsam genutzten Speicher verwenden, ohne dass tatsächlich zugewiesen wird. In Ihrem Kernel Sie haben:

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

Aber wenn man den Kernel starten Sie nicht angeben, wie viel Speicher geteilt für jeden Block zuzuordnen:

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

Die <<< >>> Syntax tatsächlich nimmt vier Argumente, wo die dritte und vierte sind optional. Die vierte ist der Stream-Index, der verwendet wird, zwischen dem Rechen- und Datenübertragung (und für die gleichzeitige Ausführung kernel) zu erhalten, aber die Überlappung drittes Argument gibt die Menge an gemeinsam genutzten Speicher pro Block. In diesem Fall nehme ich an wollen Sie speichern TileWidth * TileWidth schwimmt in dem gemeinsam genutzten Speicher, so verwenden Sie:

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

Das Hauptproblem

Wie Sie in Ihrem Kommentar erwähnen, war das eigentliche Problem, dass Ihre Matrix Breite nicht ein Vielfaches der Blockbreite war (und Höhe, da sie quadratisch ist, über das Ende der Fäden Sinn über das Ende des Arrays zugreifen würde. Die Code sollte entweder die nicht-mehrfach Fall behandeln oder es sollte sicherstellen, dass die Breite ein Vielfaches der Blockgröße ist.

Ich soll dies früher vorgeschlagen habe, aber es ist oft nützlich cuda-memcheck laufen für memeory Zugriffsverletzungen wie dies zu überprüfen.

Answer 2

Sie haben die Treiber Timeout-Einstellungen ändern zu können, ist Windows fehlerhafte Treiber zu verhindern, verfügt über das System reagiert nicht mehr zu machen. Überprüfen Sie die Microsoft beschreibt, wie das zu tun.

Answer 3

Sie sollten auch prüfen, die „timeout“ Flag Einstellung auf Ihrem GPU-Gerät. Wenn Sie das CUDA SDK installiert haben, glaube ich, die „DEVICE“ app diese Eigenschaft berichten.