CUDA Matrix descansos de multiplicación de matrices grandes

Question

Tengo el siguiente código de la multiplicación de matrices, implementado usando CUDA 3.2 y VS 2008. Estoy funcionando en el servidor Windows 2008 R2 de la empresa. Me postulo una Nvidia GTX 480. El siguiente código funciona bien con valores de "Ancho" (Matrix ancho) 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:

Me parecía en línea y vi que algunas personas tienen este problema debido a que el organismo de control estaba matando el núcleo después de que se cuelga durante más de 5 segundos. Probé la edición del "TdrDelay" en el registro y esto retrasó el tiempo antes de que la pantalla en negro y el mismo error apareció. Así que llegué a la conclusión que no era mi problema.

Me depurando en mi código y encontró esta línea para 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 asignación de memoria correctamente y no puedo entender por qué esto está ocurriendo. Pensé que tal vez yo no tenía suficiente memoria en mi tarjeta para esto, pero entonces no debería cudaMalloc he devuelto un error? (Me confirmó que no lo hizo durante la depuración).

Cualquier idea / ayuda sería muy apreciada! ... Muchas gracias chicos !!

código del núcleo:

//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 no es en realidad el núcleo del cudaMemcpy(). Al iniciar el núcleo de la GPU se apaga y hace el trabajo de forma asíncrona con la CPU, por lo que sólo cuando se sincroniza con la GPU que usted tiene que esperar a que el trabajo a fin. cudaMemcpy() implica una sincronización implícita, por lo tanto, que es donde se ve el problema.

Se podría duplicarse a comprobar esto llamando cudaThreadSynchronize() después de que el núcleo y el problema aparece estar en el lugar de la cudaThreadSynchronize() cudaMemcpy().

Después de cambiar el tiempo de espera de TDR, ¿se reinicie su máquina? Desafortunadamente Windows necesita ser reiniciado para cambiar la configuración del TDR. Este documento de Microsoft tiene una bastante buena descripción de los ajustes completos disponibles.

problemas con el núcleo

En este caso el problema no es en realidad el tiempo de espera WDDM. Hay errores en el núcleo que se necesitaría para resolver (por ejemplo, usted debería ser capaz de i incremement por más de uno en cada iteración) y la salida a la muestra matrixMul en el SDK puede ser útil. Por cierto, espero que este es un ejercicio de aprendizaje ya que en realidad sería mejor (para un rendimiento) usando CUBLAS para llevar a cabo la multiplicación de matrices.

El problema más crítico en el código es que está utilizando memoria compartida sin tener que asignar ninguna. En su núcleo tiene:

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

Sin embargo, cuando se inicia el kernel no se especifica cómo compartió la cantidad de memoria para asignar a cada bloque:

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

La sintaxis <<< >>> realmente tiene cuatro argumentos en la tercera y cuarta son opcionales. El cuarto es el índice de flujo que se utiliza para obtener superposición entre cálculo y transferencia de datos (y para la ejecución concurrente de kernels) pero el terceros argumento especifica la cantidad de memoria compartida por bloque. En este caso supongo que deseas flotadores tienda TileWidth * TileWidth 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 problema principal

Como se menciona en su comentario, el problema real era que su matriz de anchura no era un múltiplo del ancho del bloque (y la altura, ya que es cuadrado, es decir, los hilos más allá del final accedería más allá del final de la matriz. La código debe manejar, ya sea el caso no múltiple o debe asegurar que la anchura es un múltiplo del tamaño de bloque.

Me debería haber sugerido esto antes, pero a menudo es útil ejecutar cuda-memcheck para comprobar si hay violaciónes de acceso memeory como éste.

Answer 2

Hay que cambiar la configuración de tiempo de espera del conductor, es ventanas función para evitar que los conductores defectuosos para hacer que el sistema deje de responder. Compruebe el Microsoft Página que describe cómo hacerlo.

Answer 3

También debe comprobar el ajuste de la bandera de "tiempo de espera" en su dispositivo GPU. Si tiene instalado el SDK de CUDA, creo que la aplicación "deviceQuery" informará de esta propiedad.