Tópicos sobre GPGPU em CUDA Paulo A. Pagliosa [email protected] Faculdade de Computação Universidade Federal de Mato Grosso do Sul Natal, maio de

Tópicos sobre GPGPU em CUDAPaulo A. Pagliosa

[email protected]

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


• 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



5

Processamento de geometria


Transformação

Iluminação

Projeção

Montagem de primitivos

Recorte

Primitivos

Hardware

Exibição

Framebuffer

Rasterização





6

Mapeamento de textura

Aplicação de neblina

Composição

Tonalização

Hardware

Exibição

Framebuffer

Rasterização





7


• 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




• 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




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

10


• 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


• 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


• 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


14

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


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

15

http://www.amazon.com/OpenGL-Shading-Language-Randi-Rost/dp/0321637631

http://www.amazon.com/OpenGL-SuperBible-Comprehensive-Tutorial-Reference/dp/0321498828/ref=pd_sim_b_1


• 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


• 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


• 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


20

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


• 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


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

22


• 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


• 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


• 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


• 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


• 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


• 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


• 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


• 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


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


• 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


• 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


• 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


• 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


37

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


• 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


39

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


• Exemplo: transposição de matrizes

40

__syncthreads()

Mem. compartilhada


• Exemplo: transposição de matrizes

41

Mem. compartilhada


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


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.


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


• Memória global– Com coalescência

46


47

• Memória global– Sem coalescência


48

• Memória global– Sem coalescência


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


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


51

• Mem. compartilhada– Sem conflitos


• 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


• 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


• 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


• 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


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

68


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

69


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

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

70


• 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


• 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


• 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


• 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


• 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


• 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


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

78


79


80


81


82


• 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


84

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

𝑏0𝑡 0

𝑡1

𝑡5

𝑡 2

𝑡 3

𝑡 4

𝑡 6

𝑝1 𝑝2 𝑝3



85

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

𝑏0𝑏1𝑏2𝑡 0

𝑡1

𝑡5

𝑡 2

𝑡 3

𝑡 4

𝑡 6

𝑝1

𝑝2

𝑝3



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


• 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


• 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


• 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


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


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


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


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


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


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

Documents

Tópicos sobre GPGPU em CUDA Paulo A. Pagliosa [email protected] Faculdade de Computação Universidade Federal de Mato Grosso do Sul Natal, maio de