Tópicos sobre GPGPU em CUDA Paulo A. Pagliosa pagliosa@facom.ufms.br Faculdade de Computação...

Tópicos sobre GPGPU em CUDAPaulo A. Pagliosa

pagliosa@facom.ufms.br

Faculdade de ComputaçãoUniversidade Federal de Mato Grosso do Sul

Natal, maio de 2013

Apresentação

• Conteúdo– Introdução a GPGPU (www.gpgpu.org)

– O que é CUDA– Fundamentos de programação CUDA– Práticas de programação– APIs e ferramentas– Estudo de caso

2

Introdução a GPGPU

• GPUs: unidades de processamento gráfico– Originalmente projetadas para processamento

gráfico 3D– Exemplos

3

NVIDIA GeForce 8800 GTXNVIDIA Tesla C1060

Introdução a GPGPU

• Pipeline gráfico fixo (OpenGL)

4

HardwareSoftware

Aplicação Processamento de geometria

Exibição

Processamento de fragmentos

Cena

Dados de geometria

Dados de pixels

Fontes de luz

Materiais

Rasterização

Framebuffer

Hardware

Exibição

Rasterização

Framebuffer

Introdução a GPGPU

• Pipeline gráfico fixo (OpenGL)

5

Processamento de geometria

Processamento de fragmentos

Transformação

Iluminação

Projeção

Montagem de primitivos

Recorte

Primitivos

Hardware

Exibição

Framebuffer

Rasterização

Processamento de geometria

Processamento de fragmentos

Introdução a GPGPU

• Pipeline gráfico fixo (OpenGL)

6

Mapeamento de textura

Aplicação de neblina

Composição

Tonalização

Hardware

Exibição

Framebuffer

Rasterização

Processamento de geometria

Processamento de fragmentos

Introdução a GPGPU

• Pipeline gráfico fixo (OpenGL)

7

Introdução a GPGPU

• GPUs de primeira geração (1998)– Pipeline fixo– Renderização de triângulos pré-transformados– Exemplos: TNT2, Voodoo3

• GPUs de segunda geração (1999-2000)– Transformações geométricas– Iluminação– Velocidade de renderização maior– Exemplos: GeForce 2, ATI 5700

8

Hardware

Exibição

Framebuffer

Rasterização

Processamento de geometria

Processamento de fragmentos

Introdução a GPGPU

• Pipeline gráfico programável (OpenGL 3.1)

9

Transformação

Iluminação

Projeção

Montagem de primitivos

Recorte

Primitivos

Mapeamento de textura

Aplicação de neblina

Composição

Tonalização

Hardware

Exibição

Framebuffer

Rasterização

Processamento de geometria

Processamento de fragmentos

Introdução a GPGPU

• Pipeline gráfico programável (OpenGL 3.1)

10

Introdução a GPGPU

• Shaders– Módulos que executam em GPU– Um shader pode ser de:• Vértice• Geometria (SM 4)• Fragmento

– Substituem a funcionalidade do pipeline fixo• Programa em GPU: um ou mais shaders

11

Programa GPU

Shader de vértice Shader de geometria Shader de fragmento

Introdução a GPGPU

• Shader de fragmento– Opera isoladamente sobre um fragmento– Entrada

• Variáveis pré-calculadas pela OpenGL• Variáveis definidas pela aplicação

– Saída: cor do fragmento– Operações

• Mapeamento de textura• Tonalização• Aplicação de neblina

– Fragmento pode ser descartado– Coordenadas do fragmento não podem ser mudadas

12

Introdução a GPGPU

• GPUs de terceira geração (2001-2002)– Pipeline programável com shaders de vértices– Número limitado de instruções sem ponto flutuante– Programação em linguagem de montagem– Exemplos: GeForce3 e 4, ATI 8500

• GPUs de quarta geração (2003-2006)– Shaders de vértices e fragmentos– Expansão do número de instruções– Ponto flutuante de 32 bits e dados em texturas– Surgimento de linguagens de shaders – Exemplos: GeForce FX, 6 e 7; ATI 9700 e 9800

13

Introdução a GPGPU

14

• NVIDIA GeForce 6800– Até 6 processadores de vértices– 16 processadores de fragmentos

Introdução a GPGPU

• Linguagens de shaders– Cg (C for graphics), NVIDIA– HLSL (High Level Shader Language), Microsoft– GLSL (OpenGL Shader Language)

15

Introdução a GPGPU

• GPGPU: programação genérica em GPU• GPGPU com shaders– Modelo de programação: streams e kernels– Streams: dados de entrada e saída• Vetores em CPU — texturas em GPU

– Kernels: shaders (de fragmento)– Saída: renderização em textura– Execução: rasterização– Mapping Computational Concepts to GPUs (GPU Gems 2)

– Dificuldades

17

Introdução a GPGPU

• Exemplos– Simulação dinâmica de corpos rígidos

– Multiplicação de matrizes potência de dois18

•Sem restrições •Com detecção de colisões

Introdução a GPGPU

• Desvantagens de GPGPU com shaders– GPU é programada através de uma API gráfica

• Curva de aprendizado da API• Overhead para aplicações não gráficas

– Flexibilidade• Memória da GPU pode ser lida (gather) mas não pode ser

escrita (scatter) de maneira geral• Shader de fragmento produz apenas saídas RGBA

• CUDA: Compute Unified Device Architecture• GPUs de quinta geração (2007-)– Computação de propósito geral– Exemplos: GeForce 8, 9, 100 e 200

19

Introdução a GPGPU

20

• NVIDIA GeForce 8800 GTX– 16 multiprocessadores (SMs)– 8 processadores (SPs) por multiprocessador

Introdução a GPGPU

• Aplicações CUDA– Visualização e simulação

– http://www.nvidia.com/object/cuda_home_new.html 21

•Traçado de raios •Simulação de fluídos

Introdução a GPGPU

• CUDA: arquitetura de computação paralela para GPGPU

22

Introdução a GPGPU

• Por que CUDA?– Desempenho versus custo– NVIDIA é líder de mercado: mais de 108 GPUs– Programação paralela para muitos– Tesla, Fermi e Kepler

• Instalação de CUDA1. Driver2. CUDA toolkit (compilador e bibliotecas)3. CUDA SDK (utilitários e exemplos)

• https://developer.nvidia.com/cuda-downloads23

Introdução a GPGPU

• Bibliografia (http://docs.nvidia.com/cuda/index.html)

– CUDA Programming Guide– CUDA Reference Manual– CUDA Best Practices Guide

24

Fundamentos de programação CUDA

• Modelo de programação– Host: executa a aplicação (CPU)– Dispositivo (GPU)• Coprocessador da CPU• Executa kernels

– Host e dispositivo tem DRAMs próprias– Host:• Aloca memória no dispositivo• Transfere dados de entrada para o dispositivo• Dispara a execução de kernels• Transfere dados resultantes do dispositivo• Libera memória no dispositivo 25

Fundamentos de programação CUDA

• Modelo de programação– Kernel• Função geralmente escrita em C para CUDA• Executa no dispositivo N vezes em N threads em paralelo

– Threads são organizadas em blocos– Um bloco é um arranjo 1D, 2D ou 3D de threads• Cada thread de um bloco tem um índice 1D, 2D ou 3D

– Blocos são organizados em grids– Um grid é um arranjo 1D ou 2D de blocos• Cada bloco de um grid tem um índice 1D ou 2D• Os blocos de um grid têm o mesmo número de threads

26

Fundamentos de programação CUDA

• Exemplo: kernel executando em 72 threads– Grid 2D com:• Dimensão 3×2×1• 6 blocos

– Blocos 2D com:• Dimensão 4×3×1• 12 threads cada

27

Fundamentos de programação CUDA

• Modelo de programação– Um kernel é uma função que:• Começa com o especificador __global__• Tem tipo de retorno void

– Kernels podem invocar outras funções que:• São especificadas como __device__• Podem invocar outras especificadas como __device__

(GPUs Kepler podem invocar kernels)– Funções que executam no dispositivo:• Não admitem número variável de argumentos• Não admitem variáveis estáticas• GPUs não Fermi não admitem recursão nem variáveis do

tipo endereço de função 28

Fundamentos de programação CUDA

• Modelo de programação– Kernels são invocados do host (ou de Kepler)– Um dispositivo executa um kernel de cada vez

(Fermi pode mais)

29

Fundamentos de programação CUDA

• Modelo de programação– Configuração de execução de um kernel

• Dimensões do grid e dos blocos• Tamanho da memória compartilhada (opcional)• Especificada na invocação (ou lançamento) do kernel

– Dimensão 3D é representada por objeto do tipo dim3– Um grid:

• 1D de dimensão dim3 dG tem dG.x × 1 × 1 blocos• 2D de dimensão dim3 dG tem dG.x × dG.y × 1 blocos

– Um bloco:• 1D de dimensão dim3 dB tem dB.x × 1 × 1 threads• 2D de dimensão dim3 dB tem dB.x × dB.y × 1 threads• 3D de dimensão dim3 dB tem dB.x × dB.y × dB.z threads

30

Fundamentos de programação CUDA

• Modelo de programação– Identificador global de uma thread– Pode ser usado para indexar vetores e matrizes em

funções __global__ ou __device__ – Determinado a partir das variáveis pré-definidas:

• dim3 gridDim– Dimensão do grid

• dim3 blockDim– Dimensão do bloco

• dim3 blockIdx– Índice do bloco no grid

• dim3 threadIdx– Índice da thread no bloco 31

Fundamentos de programação CUDA

32

• Modelo de programação– Hierarquia de memória acessada por uma thread • Memória compartilhada do bloco da thread

– Visível para todas as threads do bloco– Tempo de vida do bloco

• Memória local da thread• Memória global• Memória constante• Memória de textura Somente leitura

Tempo de vida da aplicação

Fundamentos de programação CUDA

• Modelo de programação– Capacidade de computação• 1.x (Tesla)• 2.x (Fermi)• 3.x (Kepler)

– Especificações: depende do dispositivo• Número de multiprocessadores e processadores• Dimensões de grids e blocos• DRAM• Memória compartilhada, etc.

33

Fundamentos de programação CUDA

• Interface de programação– C para CUDA e API de runtime– API de driver– Ambas APIs têm funções para:• Gerência de memória no dispositivo• Transferência de dados entre host e dispositivo• Gerência de sistemas com vários dispositivos, etc.

– API de runtime tem também funções para:• Gerenciamento de threads• Detecção de erros

– Manual de referência de CUDA34

APIs exclusivas

Fundamentos de programação CUDA

• Programa básico em C para CUDA– Seleção do dispositivo a ser usado– Alocação de memória no host– Dados de entrada na memória do host– Alocação de memória no dispositivo– Transferência de dados do host para dispositivo– Invocação do(s) kernel(s)– Transferência de dados do dispositivo para host– Liberação de memória no host– Liberação de memória no dispositivo– Finalização

35

Fundamentos de programação CUDA

• Programa básico em C para CUDA– Seleção do dispositivo a ser usadocudaSetDevice()

– Alocação de memória no dispositivocudaMalloc()

– Transferência de dados entre host e dispositivocudaMemcpy()

– Liberação de memória no dispositivocudaFree()

– FinalizaçãocudaDeviceReset()

36

Fundamentos de programação CUDA

37

• Exemplo: multiplicação de matrizes– Sem memória compartilhada

Fundamentos de programação CUDA

• Memória compartilhada– Disponível para threads de um mesmo bloco– Acesso 400 a 600 mais rápido que memória global– Permite colaboração entre threads (do bloco)

• Cada thread do bloco transfere dados da memória global para memória compartilhada

• Após a transferência, threads devem ser sincronizadas com __syncthreads()

• Kernel efetua operações usando os dados da memória compartilhada: cada thread pode usar dados carregados por outras threads do bloco

• Se necessário, threads podem ser sincronizadas com __syncthreads(); cada thread do bloco transfere dados da memória compartilhada para memória global

38

Fundamentos de programação CUDA

39

• Exemplo: multiplicação de matrizes– Com memória compartilhada

Fundamentos de programação CUDA

• Exemplo: transposição de matrizes

40

__syncthreads()

Mem. compartilhada

Fundamentos de programação CUDA

• Exemplo: transposição de matrizes

41

Mem. compartilhada

Fundamentos de programação CUDA

42

• Implementação em hardware– Arquitetura: arranjo de multiprocessadores (SMs)– Cada SM consiste de:• 8 processadores (SPs)• 1 unidade de instrução• Memória compartilhada

– Cada SM:• Executa threads de um bloco em grupos de 32: warp• Implementa barreira de sincronização:

__syncthreads()

• Emprega arquitetura SIMT:Single Instruction Multiple Thread

Dicas de performance

43

• Estratégias básicas– Maximização da execução em paralelo

• Estruturação do algoritmo• Escolha da configuração de execução do kernel

– Número de threads por bloco múltiplo do tamanho do warp– Mínimo de 64 threads por bloco– Configuração inicial: entre 128 e 256 threads por bloco

• Evitar divergência dentro do mesmo warp

– Otimização do uso de memória• Minimização de transferência de dados host/dispositivo• Acesso coalescido à memória global• Uso de mem. compartilhada• Acesso sem conflitos de bancos à mem. compartilhada

– Otimização do uso de instruções

Dicas de performance

44

• Memória global– O dispositivo é capaz de ler palavras de 4, 8 ou 16

bytes da memória global para registradores com UMA única instrução, DESDE que o endereço de leitura seja alinhado a (múltiplo de) 4, 8, ou 16.

– A largura de banda da memória global é mais eficientemente usada quando acessos simultâneos à memória por threads de um meio-warp (durante a execução de uma instrução de leitura ou escrita) podem ser coalescidos em uma única transação de 32, 64 ou 128 bytes de memória.

Dicas de performance

45

• Memória global– Coalescência em dispositivos 1.0 e 1.1

• Ocorre em transação de 64 ou 128 ou duas transações de 128 bytes

• Threads de um meio-warp devem acessar:– Palavras de 4 bytes, resultando numa transação de 64 bytes,– Ou palavras de 8 bytes, resultando numa transação de 128 bytes,– Ou palavras de 16 bytes, resultando em duas transações de 128 bytes

• Todas as 16 palavras (cada uma acessada por uma das 16 threads do meio-warp) devem estar no mesmo segmento de tamanho igual ao tamanho das transações (ou seja, 64, 128 ou 256 bytes). Como consequência, o segmento deve ser alinhado a este tamanho

• Threads devem acessar palavras na sequência: a thread k deve acessar a palavra k

Dicas de performance

• Memória global– Com coalescência

46

Dicas de performance

47

• Memória global– Sem coalescência

Dicas de performance

48

• Memória global– Sem coalescência

Dicas de performance

49

• Memória global– Coalescência com vetor de estruturas (AOS)• Tamanho da estrutura até 16 bytes

– Alinhamento deve ser 4, 8 ou 16 bytes, dependendo do tamanho

– Elementos devem estar no mesmo segmento da transação

• Estruturas maiores que 16 bytes: reorganizar em estrutura de vetores (SOA) com elementos de tamanho até 16 bytes

Dicas de performance

50

• Memória compartilhada– Dividida em módulos chamados bancos

• Em dispositivos 1.x o número de bancos é 16• Bancos tem largura de 32 bits

– Palavras sucessivas de 32 bits estão em bancos sucessivos

– Acessos a n endereços que estão em n bancos distintos são efetuados simultaneamente

– Acessos a dois endereços que estão no mesmo banco geram um conflito de banco: acessos são serializados

– Threads de um meio-warp devem acessar endereços em bancos distintos para evitar conflitos de bancos

Dicas de performance

51

• Mem. compartilhada– Sem conflitos

Dicas de performance

• Mem. compartilhada– Com conflitos

52

APIs para CUDA

• CUBLAS• CUFFT• Primitivos paralelos: Thrust (CUDA 5.0)– Algoritmos

• Soma prefixa• Redução• Compactação• Ordenação

– Exemplos no SDK CUDA• Scan• Partículas• Etc. 53

Primitivos paralelos

• Soma prefixa (scan)– Entrada

• Sequência A de n elementos de um tipo T• Operador binário associativo op com identidade I

– Saída: sequência B de n elementos do tipo T

– Exemplo: T = int, op = +, I = 0

54

1 2 3 4 5 6 7 8

0 1 3 6 10 15 21 28

Entrada A

Saída B

Soma prefixa exclusiva

36

Primitivos paralelos

• Soma prefixa (scan)– Entrada

• Sequência de n elementos de um tipo T• Operador binário associativo op com identidade I

– Saída: sequência de n elementos do tipo T

– Exemplo: T = int, op = +, I = 0

55

1 2 3 4 5 6 7 8

1 3 6 10 15 21 28 36

Entrada A

Saída B

Soma prefixa inclusiva

Primitivos paralelos

• Redução– Entrada

• Sequência A de n elementos de um tipo T• Operador binário associativo op

– Saída: valor do tipo T igual ao último elemento da soma prefixa inclusiva da sequência de entrada A

– Exemplo: T = int, op = +

65

1 2 3 4 5 6 7 8

1 3 6 10 15 21 28 36

Entrada A

saída

Primitivos paralelos

• Compactação– Entrada

• Sequência A de n elementos de um tipo T• Sequência V de n booleanos indicando os elementos válidos

– Saída: sequência B de m elementos do tipo T, onde m ≤ n é igual à redução (soma) do vetor V

– Exemplo: T = int

66

1 0 0 1 1 0 1 0 V

1 2 3 4 5 6 7 8A 1 4 5 7B

m0 1 1 1 2 3 3 4 4Índices de B

(soma prefixa exclusiva de V)

Estudo de caso: traçado de raios

• Por que traçado de raios?– Renderização fotorealística e fisicamente correta– Mais fácil de combinar vários efeitos visuais• Sombras suaves• Iluminação indireta• Superfícies reflexivas e polidas• Transparência• Profundidade de campo• Motion blur• Etc.

67

Estudo de caso: traçado de raios

• Por que traçado de raios em GPU?– Capacidade de processamento e paralelismo– Algoritmos híbridos: RT e rasterização

68

Estudo de caso: traçado de raios

• Traçado de raios em tempo real em GPU?– Vários trabalhos na literatura– NVIDIA city demo

69

Estudo de caso: traçado de raios

• Soluções interativas NVIDIA– OptiX (www.nvidia.com/object/optix.html)

– iray (https://www.mentalimages.com/products/nvidia-iray/rendering.html)

70

Estudo de caso: traçado de raios

• Traçado de raios básico• Implementação em CUDA– Raios secundários de reflexão e transparência– Número de atores e luzes limitado à memória disponível – Traçado de raios não distribuído– Sem tratamento de textura– Geometria dos atores definida por malha de triângulos– Iluminação direta pelo modelo de Phong

71

Traçado de raios básico

• TR básico: um raio de pixel por pixel

72

Traçado de raios básico

• Raios de pixel– A partir da posição da câmera, e em direção ao centro de

cada pixel da janela de projeção, trace um raio de pixel . A cor de é a cor do pixel correspondente da imagem

– Determine os objetos da cena que intercepta– Se nenhum objeto for interceptado, então atribua à cor do

pixel a cor de fundo da cena– Senão, seja o ponto de interseção de com o objeto mais

próximo da origem de

73

Traçado de raios básico

• Raios de sombra ou luz– A partir de dispare um raio de luz em direção à cada fonte

de luz – Se não interceptar nenhum ator (opaco) da cena, então

ilumina diretamente : adicione à a contribuição da luz determinada pelo modelo de iluminação de Phong

– Dependendo das propriedades de , outros raios podem ser disparados (raio de reflexão e raio de refração)

74

Traçado de raios básico

• Raios de reflexão– Se for suficientemene polido em , um raio é disparado de

em uma direção definida pela direção de e pela normal da superfície de no ponto

– A determinação da cor de é feita executando-se recursivamente o algoritmo, que trata como um novo raio de pixel e como o ponto de vista do observador. A cor de é somada à cor de

75

Traçado de raios básico

• Raios de refração ou transparência– Se for suficientemene transparente em , um raio é

disparado de em uma direção definida pela direção de , pela normal da superfície de no ponto e pela razão dos índices de refração do meio do qual parte e de

– A determinação da cor de é feita executando-se recursivamente o algoritmo, que tratacomo um novo raio de pixel e como o ponto de vista do observador. A cor de é somada à cor de

76

Traçado de raios básico

• Condições de parada– Redução do peso de um raio em cada nível de recursão até

um peso mínimo– Nível de recursão máximo

77

Traçado de raios básico

• Exemplo– Cena com 1.384 atores 4 fontes de luz– 708.108 triângulos– 10 níveis de recursão– Imagem 1024x768 pixels

78

Traçado de raios básico

79

Traçado de raios básico

80

Traçado de raios básico

81

Traçado de raios básico

82

Traçado de raios básico

• Problemas com o traçado de raios– Aliasing

• Solução: mais raios por pixel– Super-amostragem regular ou adaptativa– Traçado de raios distribuido

– Gargalo: interseção raio/objeto• No exemplo:

– 5.377.387 raios disparados– 2.985.241 interseções raio/triângulo

• Força bruta impraticável• Solução: estruturas de aceleração

– Grade regular– Octree– Kd-tree– BVH (boundary volume hierarchy)

83

• BVH: árvore binária de AABBs

Traçado de raios básico

84

𝑡 3𝑡 4𝑡5𝑡 6𝑡 0𝑡1𝑡 2

𝑏0𝑡 0

𝑡1

𝑡5

𝑡 2

𝑡 3

𝑡 4

𝑡 6

𝑝1 𝑝2 𝑝3

• BVH: árvore binária de AABBs

Traçado de raios básico

85

𝑡 3𝑡 4𝑡1𝑡 6𝑡 0𝑡5𝑡 2

𝑏0𝑏1𝑏2𝑡 0

𝑡1

𝑡5

𝑡 2

𝑡 3

𝑡 4

𝑡 6

𝑝1

𝑝2

𝑝3

• BVH: árvore binária de AABBs

Traçado de raios básico

86

𝑡 2𝑡 4𝑡1𝑡 6𝑡 0𝑡5𝑡 3

𝑏3𝑏4𝑏0𝑏1𝑏2𝑡 0

𝑡1

𝑡5

𝑡 2

𝑡 3

𝑡 4

𝑡 6

Implementação em CUDA

• Algoritmo: megakernel– Único kernel– Uma thread por pixel: raios primários e secundários– Usa pilha de raios secundários (uma por thread)

• Cada entrada da pilha armazena origem, direção, coeficiente, e peso de um raio

• Armazenada em memória local (tão lenta quanto a global)

– Parâmetros• Ponteiros para as estruturas de dados de entrada na memória

global são agrupados em uma estrutura na memória constante• Dados da câmera virtual, luz ambiente, cor de fundo da cena e

ponteiro para a imagem a ser gerada (frame) são passados como argumentos para o kernel

87

Implementação em CUDA

• Algoritmo: megakernel1. A -ésima thread gera o raio de pixel a ser traçado e o

armazena na pilha de raios . A origem e direção do raio são determinadas em função da câmera virtual, o coeficiente do raio é a cor branca e o peso do raio é 1

2. Enquanto houver raios em , faça os passos 3 a 63. Retire o raio da pilha 4. Se não interceptar a cena, então a cor de fundo da cena é

ponderada pelo coeficiente do raio e é acumulada à cor do pixel frame[]. Vá ao passo 2

88

Implementação em CUDA

• Algoritmo: megakernel5. Trace um raio de sombrado ponto de interseção para cada

fonte de luz da cena. Se não interceptar um objeto opaco antes de atingir a fonte de luz, a contribuição de é o produto do coeficiente de pela cor calculada segundo o modelo de iluminação de Phong. Este valor é então acumulado à cor final do pixel frame[]

6. Se o material do objeto possuir propriedades reflexivas e/ou transparentes, e as condições de parada não foram atingidas, então gere o(s) raio(s) secundário(s) com base nestas propriedades e insira-o(s) em . Vá ao passo 2

89

Implementação em CUDA

• Algoritmo: megakernel– Vantagens

• É simples e funciona...• Só um kernel, só um contexto CUDA

– Problemas• Divergência do caminho de execução de threads de um warp• Acesso não coalescido à memória global• Sem uso da memória compartilhada

90

Resultados

• Cena de teste– 708.108 triângulos reflexivos– BVH com 229.037 nós– Resolução da imagem, número de luzes e nível máximo de

recursão variáveis

• Equipamento– CPU: Intel Xeon 2.4 GHz, 12 GB RAM– GPU: NVIDIA Tesla C2050 (448 núcleos, 3GB RAM)

91

Resultados

• 4 luzes, 10 níveis de recursão

92

775.938 3.105.314 4.850.202 7.946.907320x240 640x480 800x600 1024x768

0

500,000

1,000,000

1,500,000

2,000,000

2,500,000

3,000,000

3,500,000

4,000,000

Raios por segundo/Resolução

CPUGPU

Resultados

• 4 luzes, 10 níveis de recursão

93

320x240 640x480 800x600 1024x768

CPU 4.34919 17.29244 27.58072 43.57401

GPU 0.30129 0.95782 1.39151 2.11677

2.507.50

12.5017.5022.5027.5032.5037.5042.50

Tempo (s)/Resolução

CPUGPU

Resultados

• 4 luzes, 10 níveis de recursão

94

Resolução FPS Speedup Eficiência %320x240 3,3 14,4 3,2640x480 1,0 18,1 4,0800x600 0,7 19,8 4,4

1024x768 0,5 20,6 4,6

Resultados

• 1024x758 pixels, 10 níveis de recursão

95

3.54

1.17

2

5.02

8.74

2

7.94

6.90

7

13.3

85.0

33

1 2 4 8

0

1,000,000

2,000,000

3,000,000

4,000,000

5,000,000

Raios por segundo/Número de luzes

CPUGPU

Resultados

• 1024x758 pixels, 10 níveis de recursão

96

1 2 4 8

CPU 22.34137 28.406 43.57401 84.59079

GPU 1.49803 1.71694 2.11677 3.08224

5.0015.0025.0035.0045.0055.0065.0075.0085.00

Tempo (s)/Número de luzes

CPUGPU

Resultados

• 1024x758 pixels, 10 níveis de recursão

97

N. luzes FPS Speedup Eficiência %1 0,7 14,9 3,32 0,6 16,5 3,74 0,5 20,6 4,68 0,3 27,4 6,1

Resultados

• 1024x768 pixels, 4 luzes

98

3.177.096 6.008.502 7.887.115 7.946.9070 1 5 10

0

1,000,000

2,000,000

3,000,000

4,000,000

5,000,000

6,000,000

7,000,000

8,000,000

Raios por segundo/Níveis de recursão

CPUGPU

Resultados

• 1024x768 pixels, 4 luzes

99

0 1 5 10

CPU 16.77608 29.55958 43.4742 43.574

GPU 0.45668 1.07006 2.04447 2.11677

2.507.50

12.5017.5022.5027.5032.5037.5042.5047.50

Tempo (s)/Níveis de recursão

CPUGPU

Resultados

• 1024x768 pixels, 4 luzes

100

N. recursão FPS Speedup Eficiência %0 2,2 36,7 8,21 0,9 27,6 6,25 0,5 21,3 4,7

10 0,5 20,6 4,6