CUDA MATRIX Multiplication quebras para grandes matrizes
-
27-09-2019 - |
Pergunta
Eu tenho o seguinte código de multiplicação da matriz, implementado usando o CUDA 3.2 e o VS 2008. Estou executando no Windows Server 2008 R2 Enterprise. Estou executando um NVIDIA GTX 480. O código a seguir funciona bem com valores de "largura" (largura da matriz) até cerca de 2500 ou mais.
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 eu defino a "largura" para 3000 ou mais, recebo o seguinte erro após uma tela preta:
Eu olhei online e vi que algumas pessoas têm esse problema porque o cão de guarda estava matando o kernel depois que ele pendura por mais de 5 segundos. Tentei editar o "TDRDELAY" no registro e isso atrasou o tempo antes da tela preta e o mesmo erro aparecer. Então eu concluí que esse não era o meu problema.
Debugei no meu código e achei essa linha como o culpado:
err = cudaMemcpy(P, Pd, size, cudaMemcpyDeviceToHost);
É isso que eu uso para retornar meu conjunto de resultados do dispositivo depois que minha função de kernel de multiplicação de matriz é chamada. Tudo até este ponto parece correr bem. Acredito que estou alocando a memória corretamente e não consigo descobrir por que isso está acontecendo. Eu pensei que talvez não tivesse memória suficiente no meu cartão para isso, mas o Cudamalloc não deveria retornar um erro? (Confirmei que não foi depurada).
Qualquer idéia/assistência seria muito apreciada! ... muito obrigado pessoal !!
Código do 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;
}
Eu também tenho essa outra função que usa memória compartilhada e também dá o mesmo erro:
Ligar:
MatrixMultiplicationSharedMemory_Kernel<<<dimGrid, dimBlock, sizeof(float)*TileWidth*TileWidth*2>>>(Md, Nd, Pd, Width);
Código do 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;
}
Solução
Controlando o tempo limite do WDDM
O problema é na verdade o kernel não o cudaMemcpy()
. Quando você inicia o kernel, a GPU dispara e faz o trabalho de forma assíncrona com a CPU, então é somente quando você sincroniza com a GPU que precisa esperar pelo fim do trabalho. cudaMemcpy()
Envolve uma sincronização implícita, portanto, é aí que você vê o problema.
Você pode verificar novamente isso ligando cudaThreadSynchronize()
depois do kernel e o problema parecerá estar no cudaThreadSynchronize()
ao invés de cudaMemcpy()
.
Depois de alterar o tempo limite do TDR, você reiniciou sua máquina? Infelizmente, o Windows precisa ser reiniciado para alterar as configurações de TDR. Este documento da Microsoft tem uma descrição bastante boa das configurações completas disponíveis.
Problemas do kernel
Nesse caso, o problema não é realmente o tempo limite do WDDM. Existem erros no kernel que você precisaria resolver (por exemplo, você deve ser capaz de aumentar i
por mais de um em cada iteração) e verificando o matrixMul
A amostra no SDK pode ser útil. Aliás, espero que este seja um exercício de aprendizado, pois, na realidade, você estaria melhor (para desempenho) usando o Cublas para realizar a multiplicação da matriz.
O problema mais crítico do código é que você está usando a memória compartilhada sem realmente alocar nenhum. No seu kernel você tem:
//Initialize shared memory
extern __shared__ float sharedArrays[];
Mas quando você inicia o kernel, não especifica quanta memória compartilhada para alocar para cada bloco:
MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock>>>(Md, Nd, Pd, Width);
A sintaxe <<< >>> realmente leva quatro argumentos em que o terceiro e o quarto são opcionais. O quarto é o índice de fluxo que é usado para se sobrepor entre computação e transferência de dados (e para execução simultânea do kernel), mas a terceiro argumento especifica a quantidade de memória compartilhada por bloco. Nesse caso, presumo que você queira armazenar TileWidth * TileWidth
flutua na memória compartilhada, então você usaria:
MatrixMultiplicationMultiBlock_Kernel<<<dimGrid, dimBlock, dimBlock.x * dimBlock.x * sizeof(float)>>>(Md, Nd, Pd, Width);
O problema principal
Como você mencionou em seu comentário, o problema real era que a largura da sua matriz não era um múltiplo da largura do bloco (e altura, pois é quadrado, o que significa que os fios além do fim acessariam além do final da matriz. O código deve Manuseie o estojo não multiple ou deve garantir que a largura seja um múltiplo do tamanho do bloco.
Eu deveria ter sugerido isso anteriormente, mas muitas vezes é útil correr cuda-memcheck
Para verificar se há violações de acesso à memessa como essa.
Outras dicas
Você precisa alterar as configurações de tempo limite do driver, o recurso do Windows para evitar que os drivers com defeito façam o sistema não responde. Verifica a Página da Microsoft descrevendo como fazer isso.
Você também deve verificar a configuração do sinalizador "Timeout" no seu dispositivo GPU. Se você tiver o CUDA SDK instalado, acredito que o aplicativo "DeviceQuery" relatará esta propriedade.