Gpu Shared

% 15 - Otimizações em CUDA I % SuperComputação - 2019/2 % Igor Montagner, Luciano Soares

Parte 0 - Calculando desempenho de memória

Suponha que todas threads acessem a memória global para seus elementos da matriz de entrada e assuma uma GPU com as seguintes configurações

• largura de banda DRAM de 200 GB/s
• 1.500 GFLOPS (Giga Floating Operations Per Second)

Supondo que iremos acessar dados float (4 bytes) na memória global, responda:

1. Qual a largura de banda necessária para obter desempenho máximo?
2. Dada a largura de banda de 200GB/s de nossa placa, qual o desempenho máximo que pode ser obtido acessando a memória global?

Ao calcular os números acima você verá que otimizações de acesso a memória são essenciais para obter desempenho máximo ao programar em GPU. Seu trabalho nesta atividade será quantificar essas diferenças de desempenho utilizando exemplos de código para multiplicação de matrizes.

Parte 1 - multiplicação de matrizes

A multiplicação de matrizes é muito usada como exemplo em SuperComputação devido a sua ocorrência frequente em diversos tipos de problemas. Veja abaixo um possível kernel:

__global__ void MatrixMulKernel(float* M, float* N, float* P, int Width) {

  // Calcule o índice da linha do elemento em P e M
  int Row = blockIdx.y*blockDim.y+threadIdx.y;

  // Calcule o índice da coluna do elemento em P e N
  int Col = blockIdx.x*blockDim.x+threadIdx.x;

  if ((Row < Width) && (Col < Width)) {
    float Pvalue = 0;
    // cada thread calcula um elemento da sub-matriz

    for (int k = 0; k < Width; ++k) {
      Pvalue += M[Row*Width+k]*N[k*Width+Col];
    }

    P[Row*Width+Col] = Pvalue;
  }

}

O código acima acessa os dados na seguinte ordem:

O código acima só utiliza a memória global, que é lenta, e a todo momento acessa os dados na segunda matriz por coluna. Esta implementação está disponível no arquivo .....

Veja agora a implementação disponível no arquivo ..... Ela utiliza memória compartilhada para otimizar o acesso a memória e usa a técnica de tiling apresentada na aula.

Entrega

Após compreender e executar os códigos acima, compare-os com

• algoritmo sequencial em CPU
• algoritmo paralelo (usando OpenMP) em CPU

Você deverá

• escrever um parágrafo sobre as condições do seu teste dizendo qual máquina foi usada e qual tamanho de matriz.
• escolher tamanhos de matrizes "adequados" para estes testes.
• plotar um gráfico com os tempos de execução acima
• plotar um gráfico com separando, para cada teste, tempo gasto com alocação/cópia de memória e tempo gasto executando o kernel.
• escrever um parágrafo comentando o resultado de cada gráfico.