1
Tópicos em Física Computacional: Introdução a Linguagem CUDA
Carine P. Beatrici
IF – UFRGS
Aula 06: Introdução a Linguagem CUDA – Otimização do Código
2
Da Aula Passada...
● Programa que soma matrizes linearizadas;● Numero de blocos por grid:
– dim3 gridsize(ceil(float(n*n)/float(blocksize.x)),
ceil(float(n*n)/float(blocksize.y)));● Numero de threads por bloco:
– dim3 blocksize(16,16);
3
#include<stdio.h>
#include<stdlib.h>
void __global__ soma(int *a, int *b, int *s, int n);
int main(void)
{
int *a,*b, *c;
int *Ga,*Gb,*Gc;
int n, i;;
printf("\n Entre com a dimensao da matriz \n\n");
scanf("%d",&n);
dim3 blocksize(16,16);
dim3 gridsize(ceil(float(n*n)/float(blocksize.x)),ceil(float(n*n)/float(blocksize.x)));
// Alocacao de memoria para as matrizes a,b,c
a=(int *)malloc(n*n*sizeof(int));
b=(int *)malloc(n*n*sizeof(int));
c=(int *)malloc(n*n*sizeof(int));
cudaMalloc((void **)&Ga,n*n*sizeof(int ));
cudaMalloc((void **)&Gb,n*n*sizeof(int ));
cudaMalloc((void **)&Gc,n*n*sizeof(int ));
// Atribuindo valores para a e b
for (i=0;i<n*n;i++) { a[i]=i; b[i]=n*n-1;} cudaMemcpy( Ga,a,n*n*sizeof(int),cudaMemcpyHostToDevice); cudaMemcpy( Gb,b,n*n*sizeof(int),cudaMemcpyHostToDevice);
soma<<<gridsize,blocksize>>>(Ga,Gb,Gc,n);
cudaMemcpy( c, Gc, n*n*sizeof(int), cudaMemcpyDeviceToHost);
printf("c[%d]= %d a+b= %d\n",n-1,c[n*n-1],a[n*n-1]+b[n*n-1]);
free(a); free(b); free(c);
cudaFree(Ga); cudaFree(Gb); cudaFree(Gc);
}
void __global__ soma(int *a,int *b,int *s, int n)
{
int i,j;
i = blockIdx.x * blockDim.x + threadIdx.x;
j = blockIdx.y * blockDim.y + threadIdx.y;
int tid = i*n + j;
if (tid < n*n)
{
s[tid] = a[tid] + b[tid];
}
}
4
Escolha Eficiente
● Com aquela definição criamos muitos blocos desnecessários
● Consequentemente muitas threads desnecessárias;
● Podemos melhorar a definição para:– dim3 gridsize(ceil(float(n)/float(blocksize.x)),
ceil(float(n)/float(blocksize.y)));● É possível melhorar a definição do numero de
threads por bloco?
5
Escolha do Blocksize
● Número de threads por bloco depende do modelo da GPU;
● Para ver a especificação da placa usa-se o programa deviceQuery;
6
O deviceQuery
● É um programa de amostras do SDK da Nvidia;
● Mostra as características da GPU;● Pode-se usar para testar se o CUDA esta
corretamente instalado;● Para executa-lo, na linha de comando:
./deviceQuery
7
O deviceQuery
Device 0: "GeForce GTX 560" CUDA Driver Version / Runtime Version 4.2 / 4.2
CUDA Capability Major/Minor version number: 2.1
Total amount of global memory: 2048 MBytes (2147155968 bytes)
( 7) Multiprocessors x ( 48) CUDA Cores/MP: 336 CUDA Cores
GPU Clock rate: 1620 MHz (1.62 GHz)
Memory Clock rate: 2004 Mhz
Memory Bus Width: 256-bit
L2 Cache Size: 524288 bytes
Max Texture Dimension Size (x,y,z) 1D=(65536), 2D=(65536,65535), 3D=(2048,2048,2048)
Max Layered Texture Size (dim) x layers 1D=(16384) x 2048, 2D=(16384,16384) x 2048
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 49152 bytes
Total number of registers available per block: 32768
Warp size: 32
Maximum number of threads per multiprocessor: 1536
Maximum number of threads per block: 1024
8
9
10
Maximum sizes of each dimension of a block: 1024 x 1024 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 65535
Maximum memory pitch: 2147483647 bytes
Texture alignment: 512 bytes
Concurrent copy and execution: Yes with 1 copy engine(s)
Run time limit on kernels: No
Integrated GPU sharing Host Memory: No
Support host page-locked memory mapping: Yes
Concurrent kernel execution: Yes
Alignment requirement for Surfaces: Yes
Device has ECC support enabled: No
Device is using TCC driver mode: No
Device supports Unified Addressing (UVA): Yes
Device PCI Bus ID / PCI location ID: 2 / 0
11
Como Otimizar a Execução
● Em alguns casos é melhor preencher todas as threads por bloco;
● Em outros casos é melhor distribuir as threads em todo o grid, criando mais blocos com menos threads;
● Existem situações onde ter dois blocos por multiprocessor pode ser melhor, devido a troca de threads;
● É necessário estudar o problema;
12
Compilando
● nvcc -arch sm_XX nome.cu -o gpu-exec● Onde XX <= CUDA Capability visto no
deviceQuery
13
Medindo o Tempo de Execução
● Podemos fazer isso:– dentro do programa;
– Por linha de comando;
14
Medindo o Tempo de Execução
● Em C:● Inclui a biblioteca:
#include <time.h>● Declara variáveis de tempo:
clock_t tini,tfin;● Valor de inicio da contagem do tempo:
tini = clock();● Valor de término da contagem do tempo:
tfin = clock();● Intervalo de tempo em segundos:
dt = (float)(tfin - tini) / CLOCKS_PER_SEC;
15
Medindo o Tempo de Execução● Em CUDA:● Declara variáveis de tempo:
cudaEvent_t start, stop;
float time;● Valor de inicio da contagem do tempo:
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );● Valor de término da contagem do tempo:
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );● Intervalo de tempo em milissegundos:
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );● Intervalo de tempo em segundos:● time/=1000.0;
16
Medindo o Tempo de Execução
● Podemos fazer a medida de tempo de forma mais simples fora do código, em tempo de execução;
● Para programas em CPU usamos o comando time:– time ./executavel
● Para as funções da GPU temos o nvprof:– nvprof ./gpu-executavel
● Não esta instalado nos nodes da ada.
17
Regras de Programação em GPGPU
● Coloque os dados na GPGPU e os mantenha lá;
● De bastante trabalho para a GPGPU fazer;● Foque no reuso dos dados dentro da GPGPU
para evitar as limitações da banda de memoria.
18
Otimizando o Código
● Otimizar o código é a parte mais difícil do desenvolvimento de um programa CUDA.
● Hoje este processo ainda é “artesanal”, dependente do problema e da placa utilizada;
● Alguns pontos importantes a considerar são:– Divergência do controle de fluxo
– Ocupação dos processadores
– Acesso combinado (coalesced) à memória global
– Conflitos de bancos da memória compartilhada
– Chamada do Kernel
19
Divergência do Controle de Fluxo
● As thread de cada bloco são divididas em warps, contendo 16 ou 32 threads, GPUs permitem a execução simultânea de todas as threads do warp, desde que todas executem o mesmo código
● Quando threads executam códigos diferentes, dizemos que houve uma divergência na execução do código..
● Exemplos: comandos if, else, while, for, etc.
__global__ void VecAdd(float* A, float* B, float* C, int n)
{
int i =threadIdx.x;
if (i < n) C[i] = A[i] + B[i];
}
20
Ocupação dos Multiprocessadores
● O segredo para obter um bom desempenho é manter os processadores da GPU sempre ocupados (há discussões).
● Para tal: – Os blocos devem ter tamanhos múltiplos do warp;
– Usar o menor número possível de registradores por thread
● O número de blocos por multiprocessador será maior● Com mais blocos por multiprocessador, temos mais
opções de threads para execução;● Especialmente quando as threads estiverem esperando
por dados da memória global;● Melhor otimização depende do problema.
21
22
Acesso Combinado (Coalesced)
● Acesso a Matrizes por linhas ou colunas:– O principio é o mesmo;
– Se acessar a memoria em sequencia será mais rápido do que fora de sequencia;
● Quando as threads de um warp acessam a memória ao mesmo tempo, o CUDA combina os acessos em uma única requisição;
● Para tal, todas os endereços devem estar localizados em um único intervalo de 128B;
23
24
25
26
27
Acesso não sequencial
● Stride = separação entre os acessos;
● Bandwidth = banda de transmição;
29