CUDA Matrix pause di moltiplicazione per grandi matrici

Question

Ho il seguente codice di moltiplicazione di matrici, implementato utilizzando CUDA 3.2 e VS 2008. Sono in esecuzione sul server Windows 2008 R2 Enterprise. Sono in esecuzione di una Nvidia GTX 480. Il seguente codice funziona bene con i valori di "larghezza" (Matrix larghezza) fino a circa 2500 o giù di lì.

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

Quando ho creato la "larghezza" a 3000 o superiore, ottengo il seguente errore dopo una schermata nera:

I guardato online e ho visto che alcune persone ha questo problema perché il cane da guardia stava uccidendo il kernel dopo si blocca per più di 5 secondi. Ho provato la modifica del "TdrDelay" nel Registro di sistema e questo ha ritardato il tempo prima che è apparso lo schermo nero e lo stesso errore. Così ho concluso questo non era il mio problema.

I debug nel mio codice e pensano che questa linea per essere il colpevole:

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

Questo è quello che uso per restituire il mio set di risultati dal dispositivo dopo la mia funzione del kernel moltiplicazione di matrici è chiamato. Tutto ciò che fino a questo punto sembra funzionare bene. Credo che sto allocazione di memoria in modo corretto e non riesco a capire perché questo sta accadendo. Ho pensato che forse non ho avuto abbastanza memoria sulla mia carta per questo, ma poi non avrei cudaMalloc ho restituito un errore? (Mi ha confermato che non ha fatto durante il debug).

Tutte le idee / assistenza sarebbe molto apprezzato! ... Grazie ragazzi !!

Codice del kernel:

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

Ho anche questa altra funzione che utilizza la memoria dinamica, e dà anche lo stesso errore:

di chiamata:

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

Codice del kernel:

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

Controllo della Timeout WDDM

Il problema è in realtà il non kernel cudaMemcpy(). Quando si avvia il kernel la GPU si spegne e fa il lavoro in modo asincrono con la CPU, quindi è solo quando si esegue la sincronizzazione con la GPU che si deve attendere che il lavoro alla fine. cudaMemcpy() comporta una sincronizzazione implicita, quindi che è dove si vede il problema.

Si potrebbe controllare due volte questo chiamando cudaThreadSynchronize() dopo che il kernel e il problema sarà sembrano essere sulla cudaThreadSynchronize() invece del cudaMemcpy().

Dopo aver modificato il timeout TDR, hai riavviato il computer? Purtroppo Windows ha bisogno di essere riavviato per modificare le impostazioni TDR. Questo documento Microsoft ha una discreta descrizione delle impostazioni complete disponibili.

dei problemi del kernel

In questo caso il problema non è in realtà il timeout WDDM. Ci sono errori nel kernel, che si avrebbe bisogno di risolvere (ad esempio, si dovrebbe essere in grado di i incremement da più di uno su ogni iterazione) e controllando il campione matrixMul in SDK può essere utile. Per inciso, spero che questo è un esercizio di apprendimento in quanto, in realtà, si sarebbe meglio (per le prestazioni) utilizzando CUBLAS per eseguire la moltiplicazione di matrici.

Il problema più critico nel codice è che si sta utilizzando la memoria condivisa senza in realtà l'assegnazione alcuna. Nel vostro kernel si dispone di:

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

Ma quando si avvia il kernel non si specifica come memoria da allocare per ogni blocco condiviso molto:

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

Il <<< >>> sintassi in realtà prende quattro argomenti in cui la terza e la quarta sono opzionali. Il quarto è l'indice di flusso che viene utilizzato per ottenere sovrapposizione tra elaborazione e trasferimento dati (e per l'esecuzione concomitante dei kernel) ma il terzi specifica l'argomento quantità di memoria condivisa per blocco. In questo caso presumo che si desidera galleggianti negozio TileWidth * TileWidth nella memoria condivisa, in modo da usare:

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

Il problema principale

Come si menziona nel tuo commento, il vero problema è che la vostra matrice larghezza non era un multiplo della larghezza del blocco (e l'altezza dal momento che è quadrato, cioè i fili oltre la fine sarebbe accedere oltre la fine della matrice. La codice deve gestire sia il caso non multiplo o dovrebbe assicurare che la larghezza è un multiplo della dimensione del blocco.

I dovrebbe hanno suggerito questo prima, ma spesso è utile per eseguire cuda-memcheck per verificare la presenza di violazioni di accesso memeory come questo.

Answer 2

È necessario modificare le impostazioni di timeout del driver, è finestre sono dotate di prevenire driver difettosi per rendere il blocco del sistema. Controllare il Microsoft Pagina che descrive come fare.

Answer 3

Si dovrebbe anche verificare l'impostazione flag "timeout" sul vostro dispositivo GPU. Se avete installato l'SDK CUDA, credo che l'app "deviceQuery" segnalerà questa proprietà.