Início / coding / Perguntas / 76934852
Accepted
user366312
Asked: 2023-08-19 19:26:31 +0800 CST

Por que eles estão preenchendo apenas uma memória compartilhada?

  • 772

O link acima diz que:

__global__ void coalescedMultiply(float *a, float *c, int M)
{
  __shared__ float aTile[TILE_DIM][TILE_DIM],
                   transposedTile[TILE_DIM][TILE_DIM];
  int row = blockIdx.y * blockDim.y + threadIdx.y;
  int col = blockIdx.x * blockDim.x + threadIdx.x;
  float sum = 0.0f;
  aTile[threadIdx.y][threadIdx.x] = > a[row*TILE_DIM+threadIdx.x];
  transposedTile[threadIdx.x][threadIdx.y] =
      a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
      threadIdx.x];  
  __syncthreads();
  for (int i = 0; i < TILE_DIM; i++) {
      sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
  }
  c[row*M+col] = sum;
}

(...)
Esses conflitos multibanco são muito caros. A solução simples é preencher a matriz de memória compartilhada para que ela tenha uma coluna extra, como na linha de código a seguir.

__shared__ float transposedTile[TILE_DIM][TILE_DIM+1];

Minha pergunta é,

Por que eles estão preenchendo apenas a memória compartilhada do lado direito ( transposedTile), não ambos ( transposedTilee aTile)?

cuda

1 respostas

  1. Best Answer

    Pois somente o acesso a esse tile apresentará conflitos de banco, baseado nos padrões de acesso do código. Portanto, apenas esse ladrilho precisa ser preenchido para evitar conflitos de banco.

    Resumidamente, conflitos de banco em acesso compartilhado ocorrem quando 2 ou mais threads no warp estão tentando acessar valores no mesmo banco, na mesma transação (basicamente como dizer na mesma instrução emitir). Para simplificar, quando temos uma matriz de threadblock de 32x32 e um arranjo de blocos de 32x32 com quantidades de 4 bytes, bancos e colunas são sinônimos. Portanto a declaração anterior pode ser modificada para dizer que se duas ou mais threads estiverem acessando a mesma coluna no tile, no mesmo ciclo, haverá conflitos de banco.

    A exceção a esta regra é a exceção de transmissão. Se duas ou mais threads estiverem acessando o mesmo banco (ou coluna neste caso) porque estão acessando o mesmo local, então esse local será transmitido para todos os threads necessários para aquele ciclo, e a regra de broadcast não implica em conflitos de banco.

    Portanto, isso significa que devemos examinar todos os padrões de acesso (carregar e armazenar) para cada bloco e ver quais deles no código original estariam em conflito com o banco. Nossa regra é que um conflito de banco ocorrerá se a coluna for a mesma entre as threads em um warp, mas a linha for diferente. Também devemos ser capazes de observar ou perceber rapidamente o "layout" geral de um threadblock de 32x32: existem 32 warps nesse threadblock. No primeiro warp, todos os threads terão threadIdx.yvalor igual a zero, threadIdx.xvariando de 0 a 31 ao longo do warp. A única diferença para o segundo warp é que seus threadIdx.yvalores serão todos 1. E 2 para o 3º warp, e assim por diante.

    Agora vamos olhar para cada acesso (um load e um store para cada tile) e ver quais terão um padrão de banco em conflito no código original:

    aTile[threadIdx.y][threadIdx.x] = ...

    Esta é a operação de armazenamento para aTile. threadIdx.yestá no índice de linha e threadIdx.xestá no índice de coluna. Sabemos que para um determinado warp em um threadblock de 32x32, threadIdx.yé constante ao longo do warp e threadIdx.xvaria ao longo do warp de forma linear: 0,1,2,3,...31. Portanto, esta transação de carregamento compartilhado terá threads solicitando dados na mesma linha de memória compartilhada e não haverá conflitos de banco. Cada thread está lendo de uma coluna separada, o que é como dizer um banco separado neste caso, e esse é o melhor padrão canonicamente para memória compartilhada: sem conflitos de banco.

    transposedTile[threadIdx.x][threadIdx.y] = ...

    Esta é a transação de armazenamento compartilhado para transposedTile. Vemos que em relação ao caso anterior, a indexação é invertida. threadIdx.y, que é constante em um warp, agora está no índice da coluna e threadIdx.xque varia ao longo do warp, está no índice da linha. Isso significa que estamos selecionando uma única coluna no warp, com diferentes threads acessando diferentes membros dessa coluna. Este é o acesso colunar (em vez do acesso por linha no caso anterior) e quando temos um arranjo de ladrilhos de forma que cada coluna corresponda a um banco de memória compartilhada, este será canonicamente o pior padrão de acesso possível à memória compartilhada. Todas as threads estão acessando o mesmo banco, e em todo o warp teremos um conflito de banco de 32 vias, o que significa que o acesso levará 32 vezes mais do que o melhor caso possível.

    ... * transposedTile[i][threadIdx.x];

    esta é a operação load-from-shared para transposedTile. ié constante ao longo do warp (em qualquer iteração do loop) e threadIdx.xvaria ao longo do warp. Portanto, estamos selecionando uma linha específica e solicitando os elementos ao longo dessa linha. Este é o melhor caso para acesso compartilhado.

    sum += aTile[threadIdx.y][i] ...

    Considerado em todo o warp, isso tem linhas e colunas fixas/constantes em qualquer iteração de loop. Este é o caso da transmissão . Em qualquer iteração de loop e ciclo de emissão de instrução, cada thread no warp está solicitando o mesmo local na memória compartilhada. Também é canonicamente bom.

    Portanto, apenas 1 dos 4 padrões de acesso possíveis tem o potencial de ser um conflito de banco. Este padrão de acesso pertence a transposedTilee, portanto, só transposedTileprecisa ser preenchido, para evitar conflitos de banco neste código.

    • 1

relate perguntas