Racionalizando o que está acontecendo no meu kernel simples de OpenCl em relação à memória global

Question

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

O kernel acima é uma adição de vetor realizada dez vezes por loop. Eu usei o Guia de Programação e o Flack Overflow para descobrir como a memória global funciona, mas ainda não consigo descobrir o meu código se estiver acessando a memória global de uma maneira boa. Estou acessando -o de maneira contígua e acho que de uma maneira alinhada. O cartão carrega 128kb pedaços de memória global para as matrizes A, B e C? Ele carrega os pedaços de 128kb para cada matriz uma vez para cada 32 índices GID processados? (4*32 = 128) Parece que não estou desperdiçando nenhuma largura de banda de memória global, certo?

BTW, o Profiler de computação mostra uma eficiência GLD e GST de 1.00003, o que parece estranho, eu pensei que seria apenas 1,0 se todas as minhas lojas e cargas fossem coalescedas. Como está acima de 1.0?

Answer 1

Sim, seu padrão de acesso à memória é praticamente ideal. Cada Halfwarp está acessando 16 palavras consecutivas de 32 bits. Além disso, o acesso está alinhado, pois os próprios buffers estão alinhados e o StartIndex para cada meia -WARP é um múltiplo de 16. Portanto, cada Half -WARP gerará uma transação 64byte. Portanto, você não deve desperdiçar largura de banda de memória por meio de acessos descontraídos.

Como você pediu exemplos em sua última pergunta, vamos modificar esse código para outro (padrão de acesso menos ideal (já que o loop não faz nada que eu ignorarei isso):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

A princípio, permite que isso funcione no Hardware Compute 1.3 (GT200)

Para as gravações para a, isso gerará um padrão ligeiramente não ideal (seguindo os Halfwarps identificados por seu intervalo de identificação e o padrão de acesso correspondente):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

Então, basicamente, estamos desperdiçando cerca de metade da nossa largura de banda (quanto menos que a largura de acesso dobrada para o ímpar Halfwarps não ajuda muito porque gera mais acessos, o que não é mais rápido do que mais bytes, por assim dizer).

Para as leituras de B, os threads acessam apenas os elementos da matriz; portanto, para cada meio -warp, todos os acessos estão em um bloco alinhado de 128byte (o primeiro elemento está no limite de 128b, pois para esse elemento o GID é um múltiplo de 16 => O índice é um múltiplo de 32, para 4 elementos de bytes, isso significa que o deslocamento do endereço é um múltiplo de 128b). O AccessPattern se estende por todo o bloco de 128b, portanto, isso fará uma transferência de 128b para cada meia -Warp, na cintura novamente da largura de banda.

As leituras de C geram um dos pior cenários, onde cada fios índices em seu próprio bloco de 128b, de modo que cada tópico precisa de sua própria transferência, que uma mão é um pouco de um cenário de serialização (embora não seja tão ruim quanto normalmente, Como o hardware deve ser capaz de sobrepor as transferências). O que pior é o fato de que isso transferirá um bloco de 32b para cada segmento, desperdiçando 7/8 da largura de banda (acessamos 4b/thread, 32b/4b = 8, portanto, apenas 1/8 da largura de banda é utilizada). Como esse é o accessão de transposições ingênuas de matrizes, é altamente aconselhável fazer aqueles que usam memória local (falando por experiência própria).

Compute 1.0 (G80)

Aqui, o único padrão que criará um bom acesso é o original, todos os padrões no exemplo criarão acesso completamente desconhecido, desperdiçando 7/8 da largura de banda (transferência/thread 32b, veja acima). Para o hardware G80, todos os acessos onde o enésimo tópico em um meio -warp não acessa o enésimo elemento cria acessos tão desconectados

Compute 2.0 (Fermi)

Aqui, todo acesso à memória cria transações de 128b (o máximo necessário para coletar todos os dados, portanto, 16x128b no pior dos casos), no entanto, esses são armazenados em cache, tornando menos óbvio onde os dados serão transferidos. No momento, vamos supor que o cache seja grande o suficiente para conter todos os dados e não há conflitos, portanto, a cada 128B Cacheline será transferido no máximo uma vez. Vamos assumir uma execução serializada das halfwarps, por isso temos uma ocupação determinística do cache.

Os acessos a B ainda sempre transferirão blocos 128B (nenhum outro índices de encadeamento na Memoryarea corporing). O acesso ao C gerará 128b transferências por thread (pior acesso de acesso possível).

Para acessos a A, é o seguinte (tratando -os como leituras no momento):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

Portanto, para grandes matrizes, os acessos a um desperdiçarão quase sem largura de banda teoricamente. Para este exemplo, a realidade não é tão boa, já que os acessos a C destruirão o cache muito bem

Para o Profiler, presumiria que as eficiências acima de 1,0 são simplesmente resultados de imprecisões de ponto flutuante.

espero que ajude