CUDA Interrupciones de multiplicación de matrices para matrices grandes

Question

Tengo el siguiente código de multiplicación de matrices, implementado con CUDA 3.2 y VS 2008. Me estoy ejecutando en Windows Server 2008 r2 enterprise. Estoy ejecutando una Nvidia GTX 480. El siguiente código funciona bien con los valores de "Ancho" (ancho de la matriz) de hasta aproximadamente 2500 o menos.

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

Cuando me puse el "Ancho" a 3000 o mayor, me sale el siguiente error después de una pantalla en negro:

busqué en línea y vi que algunas personas tienen este problema debido a que el organismo de control estaba matando al núcleo después de que se cuelga por más de 5 segundos. Traté de editar el "TdrDelay" en el registro y esto retrasó el tiempo antes de que apareciera la pantalla en negro y el mismo error. Entonces concluí que este no era mi problema.

I depurado en mi código y encontraron esta línea de ser el culpable:

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

Esto es lo que uso para devolver mi conjunto de resultados desde el dispositivo después de mi función de multiplicación de la matriz del núcleo se llama. Todo hasta este punto parece funcionar bien. Creo que estoy asignando la memoria correctamente y no puedo entender por qué sucede esto. Pensé que tal vez no tenía suficiente memoria en mi tarjeta para esto, pero ¿no debería haber devuelto cudaMalloc un error? (Confirmé que no lo hice durante la depuración).

¡Cualquier idea/ayuda sería muy apreciada! ... Muchas gracias muchachos !! código

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

también tengo esta otra función que utiliza la memoria compartida, y también da el mismo error:

Llamar:

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

código del núcleo:

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

El control del tiempo de espera WDDM

El problema es, en realidad, no el núcleo del cudaMemcpy(). Cuando ejecutas el kernel, la GPU se apaga y realiza el trabajo de forma asíncrona con la CPU, por lo que solo cuando sincronizas con la GPU debes esperar a que el trabajo finalice. cudaMemcpy() implica una sincronización implícita, de ahí que vea el problema.

Puede verificarlo dos veces llamando al cudaThreadSynchronize() después del núcleo y aparecerá el problema en el cudaThreadSynchronize() en lugar del cudaMemcpy().

Después de cambiar el tiempo de espera de TDR, ¿reinició su máquina? Lamentablemente, Windows debe reiniciarse para cambiar la configuración de TDR. This Microsoft document tiene una descripción bastante buena de la configuración completa disponible.

problemas con el núcleo

En este caso el problema no es en realidad el tiempo de espera WDDM.Hay errores en el kernel que necesitaría resolver (por ejemplo, debería poder incrementar i en más de uno en cada iteración) y ver la muestra matrixMul en el SDK puede ser útil. A propósito, espero que este sea un ejercicio de aprendizaje ya que en realidad estarías mejor (para el rendimiento) usando CUBLAS para realizar la multiplicación de la matriz.

El problema más importante del código es que está utilizando la memoria compartida sin asignarla realmente. En su núcleo tiene:

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

Pero cuando se inicia el kernel no se especifica la cantidad de memoria compartida para asignar a cada bloque:

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

La sintaxis < < < >>> toma realidad cuatro argumentos donde el tercero y el cuarto son opcionales. El cuarto es el índice de flujo continuo que se utiliza para superponerse entre computación y transferencia de datos (y para la ejecución simultánea del kernel), pero el tercer argumento especifica la cantidad de memoria compartida por bloque. En este caso supongo que desee almacenar TileWidth * TileWidth flotadores en la memoria compartida, por lo que se debería utilizar:

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

El principal problema

Como se menciona en su comentario, el problema real era que el ancho de la matriz no era un múltiplo del ancho del bloque (y la altura ya que es cuadrado, lo que significa que los hilos más allá del extremo accederían más allá del final de la matriz. El código debería manejar el caso no múltiple o debería garantizar que el ancho es un múltiplo del tamaño del bloque.

Debería haberlo sugerido antes, pero a menudo es útil ejecutar cuda-memcheck para verificar violaciones de acceso a memeory como esta.

Answer 2

Tienes que cambiar la configuración del tiempo de espera del controlador, es la función de Windows para evitar que los controladores defectuosos hagan que el sistema no responda. Compruebe el Microsoft Page que describe cómo hacerlo.

Answer 3

También debe verificar la configuración del indicador "timeout" en su dispositivo GPU. Si tiene instalado el SDK de CUDA, creo que la aplicación "deviceQuery" informará esta propiedad.