Capítulo 11: Direções de Pesquisa em GPU sobre Escalarização e Execução Afim
Conforme descrito no Capítulo 2, as APIs de computação em GPU, como CUDA e OpenCL, apresentam um modelo de programação semelhante a MIMD que permite ao programador lançar uma grande matriz de threads escalares na GPU. Embora cada uma dessas threads escalares possa seguir seu próprio caminho de execução exclusivo e possa acessar locais de memória arbitrários, no caso comum, elas seguem um pequeno conjunto de caminhos de execução e realizam operações semelhantes.
O fluxo de controle convergente entre as threads da GPU é explorado na maioria, se não em todas, as GPUs modernas por meio do modelo de execução SIMT, onde as threads escalares são agrupadas em warps que são executados em hardware SIMD (consulte a Seção 3.1.1). Este capítulo resume uma série de pesquisas que exploram ainda mais a semelhança dessas threads escalares por meio da escalarização e da execução afim.
A principal percepção dessa pesquisa reside na observação da estrutura de valor [Kim et al., 2013] entre as threads que executam o mesmo kernel de computação. Os dois tipos de estrutura de valor, uniforme e afim, são ilustrados no kernel de computação no Exemplo 11.1.
Variável Uniforme
Uma variável que tem o mesmo valor constante para cada thread no kernel de computação. No Algoritmo 11.1, a variável a, bem como as constantes THRESHOLD e Y_MAX_VALUE, têm valores uniformes em todas as threads do kernel de computação. Uma variável uniforme pode ser armazenada em um único registro escalar e reutilizada por todas as threads no kernel de computação.
Variável Afim
Uma variável com valores que são uma função linear do ID da thread para cada thread no kernel de computação. No Algoritmo 11.1, o endereço de memória da variável y[idx] pode ser representado como uma transformação afim do ID da thread threadIdx.x:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Essa representação afim pode ser armazenada como um par de valores escalares, uma base e um passo, o que é muito mais compacto do que a expansão completa do vetor.
__global__ void vsadd( int y[], int a ) {
i
```Aqui está a tradução em português do arquivo Markdown, com os comentários traduzidos, mas o código não traduzido:
int idx = threadIdx.x; y[idx] = y[idx] + a; if ( y[idx] > THRESHOLD ) y[idx] = Y_MAX_VALUE; }
Algoritmo 11.1: Exemplo de operações escalares e afins em um kernel de computação (de [Kim et al., 2013]).
Existem várias propostas de pesquisa sobre como detectar e explorar variáveis uniformes ou afins em GPUs. O restante deste capítulo resume essas propostas nesses dois aspectos.
## Detecção de Variáveis Uniformes ou Afins
Existem duas abordagens principais para detectar a existência de variáveis uniformes ou afins em um kernel de computação GPU: Detecção Orientada por Compilador e Detecção via Hardware.
### Detecção Orientada por Compilador
Uma maneira de detectar a existência de variáveis uniformes ou afins em um kernel de computação GPU é fazê-lo por meio de uma análise especial do compilador. Isso é possível porque os modelos de programação GPU existentes, CUDA e OpenCL, já fornecem meios para o programador declarar uma variável como constante em todo o kernel de computação, bem como fornecer uma variável especial para o ID do thread. O compilador pode realizar uma análise de dependência de controle para detectar variáveis que dependem apenas de constantes e IDs de thread, e marcá-las como uniformes/afins. As operações que trabalham exclusivamente em variáveis uniformes/afins são então candidatas à escalarização.
O AMD GCN [AMD, 2012] depende do compilador para detectar variáveis uniformes e operações escalares que podem ser armazenadas e processadas por um processador escalar dedicado.
Asanovic et al. [2013] apresentam uma análise convergente e variante combinada que permite que o compilador determine as operações em um kernel de computação arbitrário que são elegíveis para escalarização e/ou transformação afim. As instruções dentro das regiões convergentes de um kernel de computação podem ser convertidas em instruções escalares/afins. Em qualquer transição de regiões divergentes para convergentes de um kernel de computação, o compilador insere uma instrução `syncwarp` para lidar com as dependências de registro induzidas pelo fluxo de controle entre as duas regiões. Asanovic et al. [2013] adotaram tAqui está a tradução em português deste arquivo Markdown. Para o código, não traduza o código, apenas traduza os comentários.
Essa análise gera operações escalares para a arquitetura Temporal-SIMT [Keckler et al., 2011, Krashinsky, 2011].
Computação Afim Desacoplada (DAC) [Wang e Lin, 2017] se baseia em uma análise de compilador semelhante para extrair candidatos escalares e afins a serem desacoplados em um warp separado. Wang e Lin [2017] aumentam o processo com uma análise afim divergente, com o objetivo de extrair sequências de instruções