LUIS HENRIQUE ALVES LOURENÇO
PROCESSAMENTO PARALELO DE ÁUDIO EM GPU
CURITIBA
2009
LUIS HENRIQUE ALVES LOURENÇO
PROCESSAMENTO PARALELO DE ÁUDIO EM GPU
Trabalho de Conclusão de Curso apresentado
como requisito parcial à obtenção do grau de
Bacharel em Ciência da Computação. Programa de Graduação, Setor de Ciências Exatas,
Universidade Federal do Paraná.
Orientador: Prof. Dr. Luis Carlos Erpen de
Bona
CURITIBA
2009
Sumário
Lista de Figuras . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
v
Resumo . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
vi
1
Introdução . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
1
2
Programação Paralela . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
4
2.1
Paralelismo em Nı́vel de Instrução . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
4
2.2
Multiprocessadores e Paralelismo em Nı́vel de Thread . . . . . . . . . . . . . . . . . . . . . . .
5
2.3
Paralelismo em Nı́vel de Dados . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
6
2.4
Compute Unified Device Architecture (CUDA) . . . . . . . . . . . . . . . . . . . . . . . . . . . .
6
2.4.1
Escondendo Processadores . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
8
2.4.2
Gerenciamento de Threads . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . .
9
2.4.3
Hierarquia de Memória . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 10
3
Áudio Digital . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 12
3.1
Processamento de Áudio Digital . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 13
3.2
Compressão de Áudio . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 13
3.3
O padrão MPEG-1 Layer III . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 14
3.3.1
Banco de Filtros Polifásicos de Análise . . . . . . . . . . . . . . . . . . . . . . . . . . . . 15
ii
iii
3.4
4
Transformação Discreta de Cosseno Modificada . . . . . . . . . . . . . . . . . . . . . 17
3.3.3
Modelagem Psicoacústica . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 18
3.3.4
Quantificação não-Uniforme . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 19
3.3.5
Codificação de Huffman . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 21
3.3.6
Formatação da Seqüência de Bits . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 21
LAME Ain’t an Mp3 Encoder . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 22
Processamento de Áudio em GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 23
4.1
Modelo de Servidor de Áudio . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 24
4.2
Um Servidor de Áudio com codificação em GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . 25
4.3
5
3.3.2
4.2.1
Detalhes da Implementação . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 27
4.2.2
Codificação em GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 27
Resultados . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 28
Conclusão . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 32
Referências Bibliográficas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 34
Anexo A -- CUDA Application Programming Interface . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 36
A.1 Extensões da Linguagem C . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 36
A.1.1 Qualificadores de Função . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 36
A.1.2 Qualificadores de Variáveis . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 37
A.1.3 Parâmetros de Configuração da Execução . . . . . . . . . . . . . . . . . . . . . . . . . . 38
A.1.4 Variáveis Pré-definidas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 39
A.1.5 O Compilador NVCC . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 39
iv
A.2 Componente de Execução Comum . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 40
A.2.1 Tipos Pré-definidos . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 41
A.2.2 Funções Matemáticas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 41
A.2.3 Funções de Tempo . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 41
A.2.4 Tipo Textura . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 42
A.3 Componente de Execução em GPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43
A.3.1 Funções Matemáticas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43
A.3.2 Função de Sincronização . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 43
A.3.3 Funções de Textura . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44
A.3.4 Funções Atômicas . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44
A.4 Componente de Execução em CPU . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 44
A.4.1 API de Execução . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 45
A.4.2 API do Driver . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 46
Anexo B -- Speaker (servidor) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 48
Anexo C -- Listener (cliente) . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 58
Anexo D -- lhal04.h . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 65
Anexo E -- psyKernel . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 67
v
Lista de Figuras
Figura 2.1 Exemplo de Soma Paralela de um Vetor
.................................
Figura 2.2 Modelo da Arquitetura NVidia (GeForce 8) [HALFHILL, 2008]
Figura 2.3 Hierarquia de Memória em CUDA
..........
6
8
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 11
Figura 3.1 Processo de Codificação MP3
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 15
Figura 3.2 Limiar Absoluto de Audição
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 19
Figura 3.3 Mascaramento de Freqüência
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 20
Figura 3.4 Quantificação não-Uniforme
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 21
Figura 4.1 Modelo de Servidor de Áudio
. . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . . 25
Resumo
Este trabalho se propõe a demonstrar que o processamento paralelo em Unidades de
Processamento Gráfico (GPU, do inglês, Graphics Processing Unit) pode ser amplamente utilizado para o processamento de áudio a fim de melhorar o desempenho dos algoritmos existentes
e permitir que mais dados sejam processados com menor latência. Com isso, pode permite-se
uma melhora sensı́vel na qualidade do conteúdo. Esse tipo de abordagem torna-se útil devido às novas tecnologias multimı́dia, como a TV digital de alta definição, os conteúdos online
(Streamming de áudio e vı́deo) e a comunicação através de meios digitais, como o VoIP ou
videoconferência.
Palavras-chave: Processamento Paralelo, GPU, CUDA, MPEG-1 Layer III, MP3, Servidor de
Áudio.
vi
1
1
Introdução
A produção e o processamento de multimı́dia no formato digital estão se popularizando
cada vez mais. Isso pode ser percebido na evolução da indústria do cinema e dos jogos, no
desenvolvimento da tv digital, nos estúdios de gravação de música, na utilização dos meios
digitais para comunicação, na popularização dos aparelhos celulares, reprodutores de vı́deo e
de música, entre outros [PEDDIE, 2001].
Um dos aspectos fundamentais para os sistemas multimı́dia é a baixa latência [LAGO,
2004], especialmente no caso de mı́dias contı́nuas como o áudio e o vı́deo. Esses sistemas
exigem baixa latência para um grande volume de dados. Uma abordagem que tem se mostrado
eficiente para aplicações que necessitam de alta capacidade de processamento ou que processam
um grande volume de dados é o processamento paralelo [HENNESSY; PATTERSON, 1990].
A exigência do mercado de processadores gráficos resultou na evolução das Unidades de Processamento Gráfico (GPU1 ) em um dispositivo altamente paralelo, com suporte a
multithreading2 , com muitos processadores de alto desempenho e com largo barramento do
memória. O grande desafio é desenvolver aplicações que permitam usar a capacidade de escalar em grau de paralelismo e, assim, aproveitar o aumento constante do número de núcleos de
processamento.
Devido a sua estrutura altamente paralela as GPUs estão deixando de ser dispositivos
exclusivos para o processamento de aplicaçõs gráficas, e começam a ser utilizadas para realizar o processamento de aplicações de propósito geral. A Programação de propósito geral em
1 em
inglês, Graphics Processing Unit
de executar vários processos simultâneamente
2 Capacidade
2
GPU (GPGPU3 ) tem como objetivo aproveitar todo o poder de processamento das GPUs que
atualmente possuem centenas de processadores independentes e diferentes tipos de memórias.
Os primeiros programas de propósito geral que aproveitavam o potencial das GPUs foram escritos através de APIs4 desenvolvidas exclusivamente para a computação gráfica, como é
o caso das bibliotecas gráficas OpenGL5 e Direct3D6 que por muito tempo foram a única forma
de criar programas capazes de utilizar as GPUs. Porém o modelo de programação voltado para
aplicações gráficas era muito confuso e não se mostrou ideal a programação de propósito geral.
Assim foram desenvolvidos modelos de programação de propósito geral para permitir que o
hardware das GPUs fosse utilizado. O modelo desenvolvido pela NVidia é o Compute Unified Device Architecture (CUDA7 ). O objetivo desse modelo é simplificar a programação em
GPU para que o programador possa se concentrar no paralelismo. Isso é possı́vel porque a API
desenvolvida abstrai o hardware da GPU enquanto bibliotecas na linguagem C simplificam o
acesso aos recursos do dispositivo. A ATI desenvolveu um modelo semelhante ao da NVidia,
o Close to Metal (CTM) [PEERCY; SEGAL; GERSTMANN, 2006]. O CTM permite que o
desenvolvedor acesse o conjunto nativo de instruções diretamente e, com isso, tenha mais flexibilidade no desenvolvimento e obtenha o melhor desempenho de seu hardware. Além dessas,
outras APIs, como a RapidMind8 que implementa abstrações das GPUs Intel e AMD, foram
criadas para possibilitar a programação de propósito geral em GPUs.
Este trabalho apresenta o estudo do modelo programação paralela CUDA e do padrão
MPEG-1 Layer III9 de compressão de áudio para demonstrar que o paralelismo das GPUs pode
ser aplicado na criação e na melhoria de técnicas que permitam melhorar o desempenho de
sistemas multimı́dia. O objetivo do trabalho é propor um modelo de Servidor de Áudio que
permita a implementação de um Servidor de Áudio que utilize o processamento em GPU. E,
3 em
ingês, General Purpose computing on Graphic Processing Units
Application Programming Inteface, é o conjunto de rotinas e padrões definidos por um software para
utilização de suas funcionalidades
5 http://www.opengl.org/
6 http://en.wikipedia.org/wiki/Direct3D
7 em inglês, Compute Unified Device Architecture
8 www.rapidmind.net
9 O padrão MPEG-1 Layer III também é conhecido como MP3
4 API,
3
com isso, possa demonstrar que a implementação de técnicas de processamento de áudio em
GPU, especialmente utilizando o modelo CUDA, é viável. O capı́tulo 2 introduz conceitos da
programação paralela necessários para em seguida entender o modelo de programação CUDA.
O capı́tulo 3 introduz os conceitos do Áudio Digital relevantes a este trabalho, apresenta um
estudo do padrão MPEG-1 Layer III e faz uma breve apresentação do codificador LAME. O
capı́tulo 4 apresenta um modelo de Servidor de Áudio e uma implementação baseada no modelo, que utiliza um codificador de áudio implementado em partes em CUDA. Além disso são
analisados os resultados dos testes realizados com o codificador e com o Servidor de Áudio
implementado. Por fim, o capı́tulo 5 conclui o trabalho.
4
2
Programação Paralela
A programação paralela em geral é relacionada com o aumento de desempenho, uma
vez que permite mais de uma execução simultaneamente. Porém esse aumento de desempenho
está associado ao grau de paralelismo de cada programa, ou seja, a capacidade de cada programa
executar suas instruções ou seu código ao mesmo tempo.
Neste capı́tulo serão abordados os tipos de paralelismo para introduzir os conceitos necessários à programação paralela e, em seguida, um estudo do modelo de programação CUDA.
Uma revisão da API do modelo CUDA encontra-se no anexo A.
2.1
Paralelismo em Nı́vel de Instrução
Desde 1985, os processadores utilizam a sobreposição da execução de instruções através
da técnica de pipelining para melhorar seu desempenho [HENNESSY; PATTERSON, 1990]. A
esta sobreposição de execuções se dá o nome de Paralelismo em Nı́vel de Instrução. Com isso é
possı́vel permitir que duas instruções seriais sejam executadas de forma paralela, uma vez que
nem toda instrução depende de sua antecessora.
A maior limitação do Paralelismo em Nı́vel de Instrução é a dependência entre as
instruções, isso significa que muitas vezes uma instrução deve esperar um ou mais estágios até
que outra instrução disponibilize o dado necessário por ela. Aumentar o paralelismo em nı́vel
de instrução significa diminuir o nı́vel de dependência entre as intruções.
5
2.2
Multiprocessadores e Paralelismo em Nı́vel de Thread
A afirmação de que os processadores convencionais estão chegando a seus limites
fı́sicos pode ser constatada pela desaceleração da melhora de desempenho destes processadores
e pela redução da melhora de desempenho trazida pelo paralelismo em nı́vel de instrução. Ou
seja, a melhora de desempenho dos processadores é cada vez menor e o paralelismo em nı́vel
de instrução melhora cada vez menos o desempenho desses dispositivos. Com isso, os multiprocessadores passam a desempenhar o papel principal na arquitetura de computadores para
continuar a melhorar o desempenho dos computadores.
Além disso a tendência por trás dos multiprocessadores é reforçada por outros fatores:
• Aumento no interesse em servidores e no desempenho dos servidores;
• Crescimento no número de aplicações de processamento intenso de dados;
• Melhora no entendimento de como usar os multiprocessadores de uma forma eficiente;
Porém, existem dois problemas: a arquitetura de multiprocessadores é um campo extenso e diverso que em sua maior parte é novo e com muitas idéias [HENNESSY; PATTERSON,
1990]. E uma grande abrangência implica necessariamente em discutir abordagens que podem
não permanecer com o tempo.
Uma Thread, ou processo leve, é uma linha de execução de um programa que possui
suas próprias variáveis de controle como o contador1 e outras estruturas, porém ela compartilha
o mesmo código e pode compartilhar a mesma região de dados com outras threads. O advento
dos multiprocessadores permite um paralelismo real no qual as threads executam em processadores diferentes simultaneamente ao invés de alternarem sua execução no mesmo processador.
Não há dependência entre threads, portanto elas podem executar de forma livre umas das outras.
1 Program
Counter (PC)
6
Figura 2.1: Exemplo de Soma Paralela de um Vetor
2.3
Paralelismo em Nı́vel de Dados
No Paralelismo em Nı́vel de Dados os dados são divididos em partes que são execu-
tadas paralelamente em unidades de processamento diferentes. O exemplo mais simples do
paralelismo em nı́vel de dados é o incremento paralelo dos valores de um vetor. Como exemplificado na Figura 2.1, que mostra a aplicação paralela de uma função que soma 4 (representada
pelos cı́rculos) ao valor de cada posição de um vetor. Dessa forma, o tempo de execução de
todas as somas equivale a execução de uma única soma, pois todas são executadas ao mesmo
tempo em unidades de processamento diferentes. As GPUs foram desenvolvidas para utilizar
esse tipo de paralelismo e permitir o processamento de grandes quantidades de dados simultaneamente.
2.4
Compute Unified Device Architecture (CUDA)
Com a compra da ATI pela AMD, a NVidia permaneceu como a maior empresa de-
senvolvedora exclusivamente de GPUs. Com isso a concorrência pelo mercado de GPUs se
concentrou nos três principais desenvolvedores: A Intel, a AMD e a NVidia. Porém, a Intel
e a AMD, como produtoras de CPUs2 , pretendem integrar núcleos gráficos aos seus processadores em um futuro breve [HALFHILL, 2008]. Essa integração pode fazer o mercado de
placas gráficas reduzir, pois a venda desse tipo de dispositivo se restringiria a aplicações que
2 Central
Processing Unit em inglês, ou Unidade Central de Processamento
7
necessitam de um desempenho gráfico realmente alto. Por outro lado, já existem placas gráficas
vendidas integradas ao computador, portanto as conseqüências da integração de núcleos gráficos
às CPUs não deve ser tão graves.
Com isso a NVidia encontrou na GPGPU3 uma forma de se diferenciar no mercado
e aproveitar ainda mais a capacidade de processamento de suas GPUs. Porém os modelos de
programação GPGPU existentes eram muito complexos. Esses modelos haviam sido criados
para o processamento gráfico (OpenGL e Direct3D). E mapear um problema de propósito geral
para o domı́nio gráfico nem sempre é uma tarefa simples. Além disso, o resultado final depende do mapeamento escolhido. Portanto, os modelos existentes se mostraram inviáveis para
a programação GPGPU. Assim, a NVidia aproveitou a oportunidade de criar um novo modelo
de programação GPGPU. E então surgiu o CUDA.
CUDA, em inglês, Compute Unified Device Architecture, é a API que implementa o
novo modelo de programação GPGPU desenvolvido pela NVidia. Seu objetivo é proporcionar
um ambiente de programação simples, por isso CUDA implementa um mecanismo de abstração
do hardware da GPU através de bibliotecas de funções nas linguagens C/C++. CUDA permite
que o programador mantenha o foco na programação paralela, pois não requer o gerenciamento
convencional de threads, esse gerenciamento é abstraı́do pela API. Além disso, a API permite que programas desenvolvidos em CUDA não deixem de funcionar devido a atualização do
harware [NVIDIA, 2008]. CUDA é uma solução para o paralelismo real em nı́vel de thread
com alto número de processadores que possui uma arquitetura amplamente difundida. Além
disso a GPU é especialmente adequada para resolver problemas que podem ser expressados por
computação de dados paralela (ou seja, o mesmo programa é executado em vários elementos de
dado paralelamente - paralelismo em nı́vel de dados) com alta intensidade aritmética (taxa de
operações aritméticas em relação a taxa de operações de memória). Como o mesmo programa é
executado para cada elemento de dado, a necessidade de fluxos de controle sofisticados é baixa.
Portanto, o modelo de programação CUDA se mostra adequado para utilizar de forma eficiente
3 Programação
de Propósito Geral em Unidades de Processamento Gráfico (GPU), em inglês, General Propose
computing on Graphics Processing Units
8
Figura 2.2: Modelo da Arquitetura NVidia (GeForce 8) [HALFHILL, 2008]
o paralelismo das GPUs.
2.4.1
Escondendo Processadores
NVidia sempre escondeu a arquitetura de suas GPUs através de uma API. Como resul-
tado disso, os programas não acessam o hardware diretamente. Ao invés disso, as funções que
manipulam diretamente o hardware estão implementadas na API.
A figura 2.2 mostra um modelo da arquitetura GPU que serviu como base para o modelo da API de programação em CUDA. As threads são executadas nos processadores de threads4 e gerenciadas pela própria arquitetura de forma transparente ao usuário. Funções CUDA,
chamadas de kernel5 , são executadas em paralelo6 por um conjunto de processadores de threads e possuem acesso a alguns tipos de memória, incluindo a memória principal da GPU e uma
memória compartilhada entre conjuntos de processadores. Além disso, CUDA permite uma
programação heterogênea, ou seja, a programação pode ser dividida entre a CPU e a GPU de
forma que o código C serial seja executado em CPU, enquanto kernels paralelos em CUDA
executam em GPU.
4 do
inglês, Thread Processors. Também são conhecidos por Stream Processors. Na arquitetura de GPU eram
conhecidos como Shaders
5 Um Kernel consiste em uma função CUDA que contém o código que será executado em GPU
6 Paralelismo em nı́vel de Thread [HENNESSY; PATTERSON, 1990]
9
2.4.2
Gerenciamento de Threads
As threads em CUDA são identificadas por blocos, ou seja, cada bloco possui um con-
junto de threads que executam o mesmo trecho de código de forma independente entre si. Os
blocos de threads podem ser unidimensionais, bidimensionais ou tridimensionais. As threads
de um mesmo bloco podem cooperar, pois têm acesso à mesma memória compartilhada. Além
disso, existem funções definidas na biblioteca CUDA que implementam barreiras [SILBERSCHATZ; GALVIN, 2000] para sincronizar a execução das threads. O número de threads por
bloco é limitado pelos recursos de memória.
Os blocos de um kernel em CUDA são identificados por grids, que podem ser unidimensionais ou bidimensionais. As threads de cada bloco em um grid executam de forma
independente das threads dos outros blocos.
A arquitetura Tesla [NVIDIA, 2008] implementada nas placas NVidia mais recentes
é baseada em um vetor de multiprocessadores7 . Quando um programa CUDA executando em
CPU chama o grid de um kernel, os blocos do grid são ordenados e distribuı́dos aos multiprocessadores da GPU. As threads de um bloco executam concorrentemente em um multiprocessador.
Ao terminar a execução de um bloco, novos blocos são lançados para ocupar os multiprocessadores vagos.
Na arquitetura Tesla, um multiprocessador consiste em 8 núcleos de processamento
escalar8 . O multiprocessador cria, gerencia e executa threads concorrentes em hardware sem
overhead de escalonamento. Ele também implementa barreiras de sincronização com uma única
instrução.
Para gerenciar centenas de threads executando diferentes programas, o multiprocessador implementa um novo tipo de arquitetura chamada de SIMT (Single Intruction, Multiple
Thread). O multiprocessador mapeia cada thread para um núcleo de processamento. E cada
thread executa de forma independente com seu próprio endereço de instrução e registradores.
7 em
8 em
inglês, Streaming Multiprocessors
inglês, Scalar Processor
10
O SIMT cria, gerencia e executa threads em grupos de até 32 threads paralelas, chamadas
warps. As threads que compõe um warp iniciam juntas no mesmo endereço, mas são livres
para executar independentemente.
Quando um multiprocessador recebe conjunto de um ou mais blocos para executar, ele
o divide em warps que serão escalonados pelo SIMT. A cada instrução, o SIMT seleciona um
warp que está pronto para executar e passa para a próxima instrução nas threads ativas. Um
warp executa uma instrução comum por vez, assim, quando todas as threads estão executando
juntas, isto é, de forma semelhante, o warp é executado com maior eficiência.
2.4.3
Hierarquia de Memória
De forma semelhante à arquitetura do SIMT, que gerência as threads, o multiproces-
sador implementa o SIMD (Single Instruction, Multiple Data), que com uma instrução simples
controla o processamento de vários elementos.
Como ilustrado na figura 2.3, cada thread pode acessar 4 tipos de memória. Cada
thread possui uma memória local privada e um conjunto de registradores de 32 bits. Cada bloco
de threads possui uma memória compartilhada9 visı́vel a todas as threads do bloco. Todas as
threads do dispositivo possuem acesso à mesma memória global10 , que é a memória principal
da GPU. E existem também as memórias de Constante11 e de Textura12 que são acessı́veis a
todas as threads. São memórias somente-leitura otimizadas utilizadas para a entrada de dados
externos ao dispositivo o que permite o acesso indireto à memória RAM da máquina13 .
Um programa pode manipular as memórias global, de contante e de textura através da
biblioteca CUDA. Isso inclui alocação de memória, liberação de memória alocada, assim como
a transferência entre a memória do computador e do dispositivo.
O número de blocos que um multiprocessador comporta depende de quantos regis9 Shared
memory
inglês, Device Memory
11 Constant memory
12 Texture memory
13 Host Memory
10 em
11
Figura 2.3: Hierarquia de Memória em CUDA
tradores por thread e quanta memória compartilhada por bloco são necessários para um dado
kernel. Se não houver registradores ou memória compartilhada suficiente por multiprocessador
para processar pelo menos um bloco, o kernel irá falhar.
Se uma instrução executada por um warp escreve na mesma posição de memória para
mais de uma thread do warp, a ordem que as escritas ocorrem é indefinida.
12
3
Áudio Digital
Um sinal analógico de áudio é um sinal elétrico que representa as vibrações mecânicas
do ar. Tais sinais possuem duas dimensões que representam a pressão do ar variando de acordo
com o tempo. Os sistemas analógicos, utilizam a voltagem do sinal elétrica para representar a
variação da pressão do ar. Porém os sistemas analógicos são bastante vulneráveis a distorções
de sinal [WATKINSON, 2001].
O áudio digital é simplesmente um meio alternativo de transportatr um sinal de áudio.
Embora existam várias maneiras de implementar isso, há um sistema conhecido por Pulse Code
Modulation (PCM [POHLMANN, 2000]), que é amplamente utilizado. No sistema PCM, o
tempo é representado de forma discreta. Dessa forma, o sinal de áudio não é composto por uma
representação contı́nua, mas por medidas em intervalos regulares. Este processo é chamado
de amostragem e a freqüência cujas amostras são medidas é chamado de taxa de amostragem.
Cada amostra ainda varia infinitamente como o sinal original, porém, assim como o tempo,
sua representação é um valor discreto. E para completar a conversão para o formato PCM,
cada amostra é representada por um valor discreto em um processo chamado quantização. Esse
processo consiste em representar a pressão do áudio no instante da captura em um valor de
amostra.
Além de não ser tão vulnerável a distorções, a representação de áudio em formato digital ofecere várias vantagens [PAN, 1993], como a reprodutibilidade, e ainda permite a aplicação
de implementações eficientes para várias funções de processamento de áudio.
13
3.1
Processamento de Áudio Digital
O processamento de áudio digital1 é empregado na gravação e armazenamento de
áudio, para mixagem de sons e produção de programas de tv, assim como em produtos comerciais como CDs. O áudio digital é, de uma forma digital, todo o caminho do microfone até
os alto-falantes, onde procesadores de sinais digitais eficientes permitem o processamento em
tempo-real. Através do processamento de áudio é possı́vel modelar o conjunto de amostras de
áudio de forma a se obter os efeitos desejados. Com o processamento de áudio é possı́vel obter
[ROADS et al., 1996]:
• Manipulação a dinâmica da amplitude do som;
• Mixagem para combinar várias faixas de áudio;
• Filtros e equalizadores para modificar o espectro de freqüência de um som;
• Efeitos de atraso (time-delay);
• Convolução, transformações simultâneas nos domı́nios do tempo e da freqüência;
• Projeção espacial, incluindo reverberação;
• Redução de ruı́do.
3.2
Compressão de Áudio
A Compressão de Áudio Digital utiliza de técnicas de processamento de áudio para
permitir o armazenamento e a transmissão de informação de áudio de forma eficiente [PAN,
1993]. Otimizar o processo de compressão de áudio permite aumentar a variedade de aplicações
para o áudio digital. Isso inclui os dispositivos de música portáteis; o áudio para cinema; rádio
e televisão digital de alta qualidade; aparelhos de DVD e muito mais [CAVAGNOLO; BIER, ].
1 Processamento
de Sinais de Áudio Digital
14
As técnicas de compressão de áudio diferem pela complexidade de seus algoritmos,
pela qualidade da compressão do áudio e pela quantidade de dados comprimidos. Técnicas
simples como a transformação µ-law2 [PAN, 1993] e a modulação diferencial adaptaviva por
códigos de pulsos (ADPCM3 ) [PAN, 1993] podem ser facilmente implementados para processar
áudio em tempo-real. O desafio é desenvolver uma implementação em tempo-real para o padrão
de áudio MPEG-1 layer III [PAN, 1995]. As próximas seções explicam alguns dos conceitos
mais importantes do padrão MPEG-1 Layer III.
3.3
O padrão MPEG-1 Layer III
O MPEG, Motion Pictures Experts Group, é o grupo formado pela ISO4 para definir
padrões de compressão e transmissão de áudio e vı́deo. Os padrões MPEG cobrem diferentes
aspectos. Dentre eles, o padrão MPEG-1, foi o primeiro a definir a codificação do áudio.
O padrão MPEG-1 de áudio efetua a compressão do áudio baseado nas limitações
fı́sicas da audição humana. O ouvido humano é capaz de detectar sons em uma faixa de
freqüência que varia de 20Hz a 20KHz. De forma que, não faz sentido armazenar todos os
dados referentes a freqüências fora dessa faixa. Além disso, dentro da faixa de freqüências
audı́veis, a percepção da audição humana obedece à uma curva (Figura 3.2) onde a percepção
da intensidade de um som varia com a freqüência. E, por fim, o ouvido humano não consegue
captar todos os sons simultaneamente, o que é conhecido como efeito de mascaramento de sons,
onde alguns sons são escondidos por outros mais fortes. Ou seja, o padrão MPEG-1 de áudio
se utiliza das limitações da percepção da audição humana para eliminar informações de áudio
sem causar alterações perceptı́veis, sendo por isso conhecido também como um algoritmo de
codificação perceptiva.
No padrão MPEG-1 Layer III, o áudio capturado no formato PCM passa por um banco
de filtros que decompõe 1152 amostras5 PCM do áudio em 32 sub-bandas de freqüências de
2 http://en.wikipedia.org/wiki/M-law
algorithm
Differential Pulse-Code Modulation
4 International Organization for Standardization
5 Um quadro MP3 é composto de 1152 amostras PCM
3 Adaptative
15
Quantificação
não-Uniforme
Sinal Digital
de Áudio(PCM)
Banco de
Filtros
FFT
MDCT
Controle de
Distorção
Controle
de Taxa
Sinal de Áudio
Codificado
Codificação de
Huffman
Formatação
da
Seqüência
de
Bits
Modelagem
Psicoacústica
Figura 3.1: Processo de Codificação MP3
mesma largura. Após esse processo, a Transformada Discreta de Cosseno Modificada [PRINCEN; BRADLEY, 1986] (MDCT6 ) é aplicada a cada amostra de cada sub-banda. Com isso, as
sub-bandas, que pertencem ao domı́nio do tempo, serão mapeadas para o domı́nio da freqüência.
Enquanto isso, aplica-se a Transformada Rápida de Fourier [DUHAMEL; VETTERLI, 1990]
(FFT7 ) nas amostras originais para revelar seu espectro sonoro. O espectro, por sua vez, passa
pela modelagem psicoacústica que determina a taxa de energia8 do sinal para o limiar de mascaramento de cada sub-banda, que será utilizada na fase de quantificação. O bloco de controle
de distorção utiliza as taxas da relação sinal / mascaramento (SMR9 ) do modelo psicoacústico
para decidir quantos bits disponibilizar para a quantificação dos sinais das sub-bandas para reduzir o ruı́do de quantificação. Em seguida as amostras quantificadas passam pela codificação
de Huffman [HUFFMAN, 1952] para reduzir a entropia das amostras. Por fim, as amostras
codificadas e suas informações são empacotadas. As subseções a seguir descrevem os detalhes
das principais operações realizadas. O processo descrito acima está ilustrado na Figura 3.1.
3.3.1
Banco de Filtros Polifásicos de Análise
O primeiro passo do processo de codificação é a passagem do sinal de áudio PCM por
um banco de filtros. O Banco de Filtros Polifásicos de Análise tem como objetivo decompor o
sinal em 32 sub-bandas. Essa decomposição agrupa as amostras de sub-banda da mesma forma
6 Modified
Discret Cosine Transform
Fourier Transform
8 Nı́vel de pressão do ar determinada pelo sinal, em decibéis (dB)
9 Signal-to-Mask Ratio
7 Fast
16
que no sinal original, porém pode causar algumas distorções10 .
A seqüência de 1152 amostras PCM de áudio de um quadro MP3 são filtradas de
maneira que cada sub-banda possua 36 amostras. O resultado do banco de filtros é definida pela
seguinte equação:
63
S[i] =
7
∑ ∑ M[i][k] ∗ (C[k + 64 j] ∗ X[k + 64 j])
k=0 j=0
Onde i é o ı́ndice de cada uma das 32 sub-bandas; S[i] é a amostra resultante para a sub-banda i e
tempo11 t; C[n] é um dos 512 coeficientes da janela de análise definida pelo padrão [PRINCEN;
BRADLEY, 1986]; X[n] é uma amostra de áudio de um buffer de amostras de 512 posições;
M[i][k] é a matriz de coeficientes da análise que é definido pela equação:
M[i][k] = cos[
(2 ∗ i + 1) ∗ (k − 16) ∗ π
]
64
Esse conjunto de equações está otimizado para reduzir o número de cálculos. Para
melhorar o entendimento desse cálculo, essas equações podem ser simplificadas na seguinte
equação de convolução:
511
St[i] =
∑ X[t − n] ∗ Hi[n]
n=0
Onde X[t] representa uma amostra de áudio no tempo t. H[i] que é definido pela
equação H[i] = h[i] ∗ cos[ (2∗i+1)∗(n−16)∗π
], representa o filtro responsável por decompor o áudio
64
em sub-bandas de freqüência adequadas, por isso é chamado de filtro polifásico. E, por fim,
h[n] é definido por −C[n] se o resultado de n/64 for ı́mpar, caso contrário h[n] = C[n].
10 Aliasing
11 O
[VAIDYANATHAN, 1987]
tempo t é representado por um inteiro múltiplo de 32 intervalos de amostra
17
3.3.2
Transformação Discreta de Cosseno Modificada
Nesse processo, as amostras das 32 sub-bandas recebidas do banco de filtros são ma-
peadas em uma transformação discreta de cosseno modificada (MDCT). Como resultado, as
amotras, que pertencem ao domı́nio do tempo, serão mapeadas no domı́nio da freqüência.
Antes de computar a MDCT, quatro funções janela12 são aplicadas às amostras. Funções
janela, são funções utilizadas em processamento de sinais para melhorar a eficiência da análise
do espectro de onda (espectro sonoro no caso do processamento de áudio). O padrão MPEG-1
Layer III especifica dois tamanhos de blocos MDCT: o bloco longo de 18 amostras e o bloco
curto de 6 amostras. Há 50% de sobreposição entre sucessivas janelas de transformadas uma
vez que o tamanho da janela é 36 ou 12 respectivamente. Da mesma forma, dependendo da
dinâmica de cada sub-banda são usadas janelas longas ou curtas. Se as amostras de uma dada
sub-banda se comportam de forma estacionária, a janela regular, longa, é usada. Se as amostras
são transitórias, a janela curta é aplicada para subdividir o resultado da sub-banda em freqüência
e intensificar a resolução de tempo. Este mecanismo ajuda a evitar o aparecimento do fenômeno
de pré-eco13 , o que pode acontecer quando aplicamos a FFT sobre um conjunto de amostras.
O efeito de pré-eco ocorre quando há uma demanda muito alta de bits em um curto espaço de
tempo (por exemplo, um momento de silêncio seguido de um ataque abrupto), e com isso o
ruı́do de quantização exagerado de um determinado trecho de áudio é espalhado para instantes
anteriores à sua ocorrência causando um ruı́do audı́vel no sinal codificado, nos instantes anteriores à ocorrência do ataque. As outras duas janelas utilizadas para manipular as transições de
longo para curto ou de curto para longo são chamadas de janela de inı́cio e janela de parada,
respectivamente. O bloco curto é um terço do bloco longo, de forma que três blocos curtos
substituem um bloco longo. O número de amostras de um quadro de amostras não é alterado
pela tamanho do bloco. Para um dado quadro de amostras, a MDCT possui 3 modos de blocos, 2 modos com o mesmo tamanho de blocos (longos ou curtos) e um modo misto, onde as
duas sub-bandas de mais baixa freqüência usam blocos de longos e as 30 sub-bandas de mais
12 http://en.wikipedia.org/wiki/Window
function
13 http://wiki.hydrogenaudio.org/index.php?title=Pre
echo
18
alta freqüência utilizam blocos curtos. Assim é possı́vel fornecer melhor resolução para as
freqüências mais baixas, sem sacrificar as resolução de tempo para as freqüências mais altas.
3.3.3
Modelagem Psicoacústica
A modelagem psicoacústica é o componente chave para o desempenho do codificador.
Com ela é possı́vel simular a percepção do som pelo sistema auditivo humano. Na codificação,
a modelagem psicoacústica decide quais partes são acusticamente irrelevantes e quais não são,
e remove as partes inaudı́veis. Para isso, ela se aproveita da falta habilidade do sistema auditivo humano em ouvir sons quantificados sobre um mascaramento. O mascaramento é uma
propriedade do sistema auditivo humano que ocorre quando um sinal de áudio forte se encontra
próximo de um sinal de áudio mais fraco no espectro ou no tempo, tornando o sinal de áudio
mais fraco imperceptı́vel.
O limiar absoluto de audição14 consiste na quantidade mı́nima de energia necessária
para um tom puro ser detectado em um ambiente silencioso. Se a energia de um número de tons
de freqüência for medida, obtêm-se o gráfico da figura 3.2. Isso significa que todos os valores
abaixo da linha do gráfico não podem ser detectados.
O mascaramento de freqüência é um fenômeno que torna um sinal de baixa freqüência
inaudı́vel pela ocorrência simultânea de um sinal mais forte em uma freqüência suficientemente
próxima. O limiar do mascaramento de freqüência pode ser medido e qualquer sinal abaixo
dele não será audı́vel, como demonstra o gráfico da Figura 3.3. O limiar depende da intensidade
sonora e da freqüência da máscara. Com isso é possı́vel intensificar o ruı́do de quantificação de
uma sub-banda o que significa que menos bits serão necesários para representar o sinal nessa
sub-banda.
Além do mascaramento no domı́nio da freqüência, também existe o mascaramento
temporal. Isso acontece quando dois sons aparecem em um intervalo muito pequeno de tempo.
O som mais forte pode mascarar o mais fraco. Os efeitos do mascaramento temporal aconte14 Absotute
Threshold of Hearing(ATH)
19
Figura 3.2: Limiar Absoluto de Audição
cem antes e depois de um som forte. Um som pode sofrer de pós-mascaramento, quando isso
acontece após um som mais forte, ou pré-mascaramento, quando o som mais forte ocorre logo
em seguida. O pré-mascaramento pode prevenir a ocorrência de pré-eco.
3.3.4
Quantificação não-Uniforme
O bloco de Quantificação não-Uniforme recebe o resultado da MDCT, uma janela de
mudança e informações de mascaramento da modelagem psicoacústica para efetuar a quantificação.
O resultado é um dado codificado de acordo com as limitações da audição humana. A Quantificação
não-Uniforme é a parte que mais consome tempo no algoritmo de codificação. Ela é dividida em
três nı́veis: A execução do bloco de Quantificação não-Uniforme que executa o loop externo,
responsável pela análise da distorção, que, por sua vez, executa o loop interno que é responsável
pela quantificação e codificação.
Como demonstrado na Figura 3.4, as amostras das sub-bandas são quantificadas em um
processo iterativo. O loop interno quantifica a entrada e incrementa o passo do quantificador até
que os dados possam ser codificados com um certo número de bits. Após a execução do loop
Nível de Pressão do Ar
20
Máscara
Limiar de audição
modificado
Freqüência
Sinal não
Mascarado
Sinais
Mascarados
Limiar de
Audição
Figura 3.3: Mascaramento de Freqüência
interno, o loop externo faz a verificação de cada fator de escala da sub-banda, se a distorção
permitida for excedida, o fator de escala é incrementado e o loop interno é executado novamente.
O loop externo, também conhecido como loop de controle de distorção, controla o
ruı́do produzido pela quantificação no loop interno. O ruı́do é eliminado pela multiplicação das
amostras por um fator de escala. O loop externo é executado até que o ruı́do permaneça abaixo
do limiar de mascaramento para cada fator de escala da sub-banda.
O loop interno, ou loop de controle de taxa, realiza a quantificação do áudio no domı́nio
da freqüência e o prepara a operação de formatação. A tabela do código de Huffman atribui palavras menores aos menores valores quantificados. O número total de bits resultados da
codificação pode exceder o número de bits disponı́veis em um quadro, isso pode ser corrigido
ajustando o ganho global para resultar em um passo de quantificação maior e, conseqüentemente, um valor quantificado menor. Essa operação é repetida com diferentes tamanhos de
passos de quantificação até que o número de bits necessários pela codificação de Huffman seja
suficientemente pequeno.
21
Loop Externo
Ajusta o fator de escala e volta ao loop interno,
Repete enquanto o ruído da quantificação não for aceitável
Amostras
de
Sub-banda
Loop Interno
Compara o ruído de
Ajusta o ganho até que o valor
quantificação de cada
quantificado seja menor que o
sub-banda com o limiar
bitrate
Áudio
de mascaramento
Codifica- bitrate
QuantiCodificado
S
S
ção de
ficação
Huffman
Ajuste
Global de
Ganho
Ajuste do
Fator
de escala
Figura 3.4: Quantificação não-Uniforme
3.3.5
Codificação de Huffman
Dependendo da implementação, a codificação de Huffman baseada em 32 tabelas
estáticas de Huffman é efetuada durante ou após a quantificação. A codificação de Huffman
fornece uma compressão sem perda de dados, portanto é capaz de reduzir o tamanho sem perda
de qualidade. Na Codificação de Huffman a entropia é baseada na distribuição estatı́stica de um
grupo de valores. Uma tabela de substituição cobrindo todos os valores é estabelecida a partir
dos dados estatı́sticos. Nessa tabela, os valores com maiores possibilidades de aparecerem nos
dados são associados a uma palavra menor e dados que raramente aparecem são associados a
palavras maiores. Entretanto, a codificação de Huffman é um código de tamanho variável e
portanto a construção da tabela de códigos não é uma tarefa trivial. As amostras são ordenadas
pela freqüência e então divididas em três faixas distintas. Isso permite que cada faixa seja codificada com um conjunto diferente de tabelas especı́ficamente ajustadas para as estatı́sticas de
cada faixa.
3.3.6
Formatação da Seqüência de Bits
A última parte da codificação consiste na produção da seqüência de bits compatı́vel
com o padrão MPEG-1 Layer III. A seqüência de bits é particionada em quadros que represen-
22
tam 1152 amostras PCM. O cabeçalho descreve a taxa de bits e a freqüência de amostragem
usadas para o audio codificado. Informações como tipo de bloco, tabelas de huffman, ganho
de sub-banda e fatores de sub-banda são selecionados. Uma técnica utilizada para ajustar a
variação do tempo de codificação é a utilização de um reservatório de bits. O codificador pode
doar alguns bits quando ele precisa de menos do que a média de bits necessária para codificar
um quadro. Em seguida, quando o codificador precisar de mais bits, ele pode emprestar do
reservatório. O codificador pode emprestar apenas bits doados de quadros passados, não pode
emprestar de quadros futuros.
3.4
LAME Ain’t an Mp3 Encoder
Considerado um dos melhores, se não o melhor, codificador MP3, o LAME 15 (acrônimo
recursivo de LAME Ain’t an Mp3 Encoder) iniciou em 1998 como um projeto open source que
visava melhorar o modelo psicoacústico, a eliminação de ruı́do e o desempenho do codificador
dist1016 da ISO, que é a implementação do padrão MPEG-1 Layer III. Portanto, a princı́pio,
o LAME não era tecnicamente um codificador (por isso o seu nome), apenas uma tentativa de
melhorar o codificador da ISO. Para evitar problemas legais, o LAME foi desenvolvido sob
uma licensa aberta (a LGPL) ao contrário do código da ISO e da patente da organização de
pesquisa alemã Fraunhofer-Gesellschaft17 , que desenvolveu o algoritmo de compressão MP3.
Em 1999, o projeto apresentou seu próprio modelo psicoacústico, chamado de GPSYCHO, que
tem como objetivo melhorar o modelo da ISO. Finalmente, em Maio de 2000, todo o código
da ISO havia sido reescrito e o LAME surgiu com seu próprio codificador desenvolvido pela
equipe de programadores open source por trás do projeto LAME.
15 lame.sourceforge.net
16 padrão
ISO 11172-3
17 http://www.fraunhofer.de
23
4
Processamento de Áudio em GPU
É possı́vel melhorar a eficiência das técnicas de processamento de áudio, como a compressão de áudio digital, através do desenvolvimento do hardware, do desenvolvimento de novas
técnicas de processamento de sinais ou de melhorias nas técnicas existentes. Nesse sentido a
programação de propósito geral em GPU pode ser vista como um meio para tornar algoritmos
de processamento de áudio digital mais eficientes através do paralelismo desses dispositivos.
Dessa forma o hardware das GPUs pode ser usado para incrementar o desempenho dos algoritmos apenas com algumas modificações nas técnicas de processamento de áudio existentes e
sem a necessidade de desenvolver novas técnicas de processamento de sinais.
Neste trabalho propomos a implementação de Servidor de Áudio que utiliza a programação
paralela em GPUs do modelo CUDA para melhorar o desempenho do processamento do áudio.
Para demonstrar que essa abordagem é possı́vel, as próximas seções deste capı́tulo descrevem
um modelo de servidor de áudio que captura o áudio de um microfone, processa o áudio e
transmite o áudio processado a um cliente que, após conectado ao servidor, recebe o áudio,
realiza um novo processamento, se necessário, e reproduz o áudio. Também é descrita a
implementação de um servidor de áudio baseado nesse modelo. O processamento de áudio
realizado por essa imlpementação é a codificação do áudio capturado. Para melhorar o desempenho da codificação, o codificador foi modificado a fim de ter sua execução paralelizada
através da implementação de funções em CUDA.
24
4.1
Modelo de Servidor de Áudio
Este trabalho começou como um estudo do modelo CUDA de programação paralela em
GPU e uma análise das possı́veis aplicações para tal tecnologia. Entre as possı́veis aplicações
para o processamento paralelo em GPU citadas estão: algoritmos genéticos, criptografia [MANAVSKI., 2007], compactação, processamento de imagens, processamento de áudio, reconhecimento de fala [CHONG et al., 2008] e VoIP. Por fim, foi decidido que este trabalho deveria
desenvolver um modelo de Servidor de Áudio e aplicar o processamento em GPU. Dessa forma,
foi possı́vel unir as propostas de trabalhar com redes de computadores, processamento paralelo
em GPU e processamento de áudio.
O modelo proposto consiste em um servidor que captura o sinal analógico do áudio
de um microfone e converte-o para o formato digital através de uma biblioteca de áudio, representada na Figura 4.1 pela Captura de Áudio. Após a captura, o áudio, já no formato digital, é
processado utilizando uma técnica de processamento de áudio implementado em CUDA. Essa
etapa corresponde ao Processamento de Áudio na Figura 4.1. Em seguida, o áudio processado
é enviado para um cliente através de um protocolo de rede. Essa comunicação entre o servidor
e o cliente está representada na Figura 4.1 pela Transferência do Áudio Processado. O cliente,
então, recebe os dados do servidor e realiza um novo processamento, caso seja necessário. E,
por fim, o áudio digital é convertido em um sinal analógico para ser executado pelas caixas de
som do cliente. Essa ação é representada na Figura 4.1 pela Reprodução do Áudio.
Esse modelo de Servidor de Áudio serve como base para demonstrar a utilização do
modelo CUDA de programação paralela em GPU para realizar o processamento de áudio. E
dessa forma mostrar que a programação paralela em GPU pode ser utilizada como um meio de
melhorar o desempenho de técnicas de processamento de áudio existentes, como por exemplo,
a codificação de áudio.
25
Processamento
do
Áudio
Envio
Processado
Processamento
do
Áudio
Figura 4.1: Modelo de Servidor de Áudio
4.2
Um Servidor de Áudio com codificação em GPU
A primeira implementação do modelo desenvolveu um Servidor de Áudio que utiliza
a biblioteca ALSA1 para capturar o sinal analógico de áudio do microfone e convertê-lo no
formato PCM de 16 bits little-endian sinalizado; ou seja, cada amostra de áudio possui valores
entre -32768 e 32767; em apenas um canal (mono); com taxa de amostragem de 44100Hz, isso
significa que são capturadas 44100 amostras de áudio por segundo. E, além disso, o Servidor
de áudio possui um socket2 configurado com o protocolo UDP [POSTEL, 1980] para enviar
os dados do áudio para um cliente. Foi implementado também um cliente que se conecta ao
servidor através do mesmo protocolo de rede, recebe o áudio do servidor e reproduz o áudio
digital utilizando o ALSA. O algoritmo do servidor consiste em capturar e enviar o áudio para o
cliente, esses procedimentos são executados repetidamente nessa ordem de até que a execução
do servidor seja interrompida pelo usuário. Da mesma forma o cliente recebe e reproduz o áudio
do servidor, esses procedimentos também são executados nessa ordem até que o servidor pare
1 Advanced
Linux Sound Architecture
2 http://en.wikipedia.org/wiki/Internet
socket
26
de enviar mais dados, pois a função de recebimento de dados do cliente é bloqueante. Com essa
implementação é possı́vel capturar o áudio em formato digital e reproduzı́-lo em outra máquina.
Porém essa implementação não representava o modelo proposto, pois não realizava nenhum processamento de áudio entre a captura e a reprodução do áudio. Para que a implementação
correspondesse ao modelo era necessário a implementação alguma técnica para processar o
áudio do Servidor de Áudio implementado. Depois de analisar algumas alternativas, como a
criptografia do áudio, foi decidido que o processamento seria uma codificação do áudio capturado. Dessa forma, foi possı́vel reduzir a quantidade de dados enviados para o cliente e,
portanto, melhorar a qualidade do áudio capturado sem aumentar o tempo de envio do áudio
para o cliente.
Então, foram pesquisados alguns codificadores de áudio, entre eles o Speex3 , que é um
codificador de áudio especializado em fala4 . Porém o codificador escolhido foi o LAME. Pois,
dessa forma, foi possı́vel aproveitar a documentação do concurso da NVidia de implementação
do codificador LAME em CUDA5 . Na página do concurso, a NVidia disponibiliza uma versão
pré-modificada da versão 3.97 do LAME que foi utilizada como base para o codificador deste
trabalho. O codificador LAME recebe o áudio em formato digital, realiza a eliminação das
partes inaudı́veis pelo ouvido humano e comprime o áudio restante gerando dados no formato
MP3.
Com isso, a API da biblioteca do codificador LAME foi utilizada para implementar
a codificação do Servidor de Áudio que codifica os dados capturados pelo microfone com a
finalidade de diminuir o tempo de envio. Enquando a implementação do LAME em CUDA tem
o objetivo de reduzir o tempo da compressão dos dados.
3 http://speex.org/
4 em
inglês, speech codec
5 http://cudacontest.nvidia.com/index.cfm?action=contest.contest&contestid=2
27
4.2.1
Detalhes da Implementação
O cliente e o servidor foram implementados na linguagem C e para as modificações
do código da biblioteca LAME foram utilizados C e CUDA. A primeira coisa a ser feita em um
ciclo é a configuração do LAME. A cada codificação o LAME deve ser configurado pois algumas variáveis de configuração não podem ser reutilizadas. A geração do áudio é feita através
de um microfone configurado pela biblioteca de áudio ALSA que captura o áudio no formato
PCM de 16 bits little-endian sinalizado, em apenas um canal, com uma taxa de amostragem
de 44100Hz. Para conter o áudio capturado, utiliza-se um buffer implementado em um vetor
de elementos do tipo short int (16 bits) com 1152 posições6 . Logo, cada leitura do microfone
captura 1152 amostras PCM.
Após a captura, o áudio é copiado para a memória da GPU e codificado utilizando
uma função de codificação da biblioteca LAME mantendo as mesmas configurações supracitadas. Os dados do áudio codificado retornados pela função são armazenados em um buffer de
unsigned char de tamanho definido através de uma função da biblioteca LAME que calcula o
tamanho do buffer de acordo com as configurações definidas.
Por fim, o buffer MP3 é enviado a um cliente em outra máquina através de um socket que utiliza o protocolo UDP. Do outro lado, o cliente recebe o buffer MP3 e realiza sua
descodificação utilizando uma função de descodificação da biblioteca LAME que retorna o
áudio em amostras PCM e na seqüência esse áudio é reproduzido através da biblioteca de áudio
ALSA.
4.2.2
Codificação em GPU
Para codificar o áudio foram utilizadas funções da biblioteca LAME, foi utilizada a
versão 3.97 do LAME disponibilizada pelo concurso da NVidia de implementação do codificador LAME em CUDA. Essa versão do LAME possui uma implementação em CUDA para
o filtro passa-alta do modelo psicoacústico e do ajuste de escala de amostra da função de
6 Tamanho
do quadro MP3
28
codificação.
O filtro passa-alta do modelo psicoacústico é usado para a detecção de ataques e,
com isso, evitar o pré-eco. A implementação paralela do filtro passa-alta permite que todas
as 576 sub-bandas executem o filtro simultaneamente. O ajuste de escala das amostras ocorre
na preparação do processo de codificação e consiste em multiplicar o valor de cada amostra PCM por um valor de escala. Este ajuste tem como objetivo alterar o volume sonoro das
amostras. A implementação dessa função em CUDA permite a execução das multiplicações
em paralelo. Essas funções foram implementadas pela NVidia no código disponibilizado para
o concurso de implementação do codificador LAME em CUDA. É importante para que estas
implementação possam ter ganhos de desempenho que o buffer de amostras seja copiado para a
memória da GPU na preparação para a codificação, pois com isso o tempo de acesso à memória
é otimizado. Portanto foi necessário incluir no Servidor de Áudio funções da API CUDA para
implementar a cópia dos dados para a GPU.
4.3
Resultados
R
As máquinas utilizadas para os testes dessa implementação possuem processador Intel
CoreTM 2 Quad 2.4GHz de 64 bits, 2GB de memória RAM, equipada com uma GPU GeForce 8600 GT com 256MB de memória global. O sistema operacional utilizado é o Debian
GNU/Linux versão 5.0 (lenny/sid) com o kernel 2.6.24-1-amd64.
Para medir o desempenho do codificador LAME com modificações implementadas
em CUDA, foi realizado um teste comparando o tempo de codificação do LAME modificado
com a versão 3.97 original do LAME7 . No teste realizado, cada implementação codificou um
conjunto de 6 arquivos no formato WAV [BORN, 1995] com diferentes tamanhos, onde 3 desses
arquivos possuiam amostras em Mono (um canal) e os outros 3 arquivos possuiam amostras em
Stereo (dois canais). Os arquivos foram codificados utilizando a configuração padrão do LAME
com taxa de bits constante8 , mantendo a escala (volume sonoro) e o número de canais. Cada
7 http://sourceforge.net/project/showfiles.php?group
8 Constant
Bit Rate (CBR)
id=290&package id=309
29
implementação codificou cada um dos arquivos 15 vezes para fins estatı́sticos, pois, enquanto
o desvio padrão do codificador LAME original (com processamento em CPU) variou entre
0,5% e 1,1%, o codificador modificado (com parte do processamento em GPU) variou entre
1,6% e 5,4%. Com esses valores, foi produzida uma média dos tempos de execução de cada
implementação do LAME para cada arquivo. Os valores estão demonstrados na Tabela 4.1.
Tamanho dos Arquivos
Tempo Médio de codificação do LAME em GPU
Maior Tempo de codificação do LAME em GPU
Menor Tempo de codificação do LAME em GPU
Tempo Médio de codificação do LAME original
Maior Tempo de codificação do LAME original
Menor Tempo de codificação do LAME original
Ganho Médio de Desemenho do LAME em GPU
3MB
0,76s
0,79s
0,76s
0,85s
0,86s
0,85s
1,12x
6MB
1,18s
1,19s
1,18s
2,06s
2,07s
2,06s
1,75x
12MB
2,48s
2,57s
2,47s
3,06s
3,07s
3,06s
1,23x
24MB
3,52s
3,69s
3,49s
6,65s
6,67s
6,64s
1,89x
41MB
8,56s
8,64s
8,50s
11,98s
12,01s
11,97s
1,41x
82MB
13,81s
13,88s
13,70s
27,08s
27,24s
27,01s
1,97x
Tabela 4.1: Tempo de execução por tamanho dos arquivos
Os valores da Tabela 4.1 mostram uma variação muito maior nos tempos do LAME
modificado (com parte do processamento em GPU) ao contrário do que acontece com os tempos
do LAME original (com processamento em CPU). De forma que para um dos testes do LAME
modificado o desvio padrão chegou ao valor de 5,4% enquanto o desvio padrão do LAME original não passou de 1,1%. Apesar disso, os valores médios dos tempos de codificação do LAME
modificado ficaram mais próximos do menor tempo. Demonstrando que os maiores tempos
ocorreram em menor quantidade no teste. Também foi possı́vel constatar que o desempenho
das implementações que utilizaram CUDA foi até 1,97 vezes maior se comparado com codificadores que não utilizaram o modelo de programação paralela. E que mesmo no pior caso,
as implementações que utilizaram CUDA foram pelo menos 1,12 vezes mais rápidas. Com
tudo isso, foi possı́vel comprovar que é possı́vel melhorar o desempenho dos algoritmos de
processamento de áudio ao utilizar um modelo de programação paralela em GPU.
Outro teste realizado foi o de desempenho do Servidor de Áudio implementado a partir
do modelo proposto neste trabalho. Esse teste tem como objetivo demonstrar que a aplicação
desenvolvida foi capaz de se utilizar plenamente do ganho de desempenho proporcionado pela
implementação de tecnicas de processamento de áudio em GPU.
30
Servidor de Áudio sem codificação
Servidor de Áudio com codificação em CPU
Servidor de Áudio com codificação em GPU
TCL
3263µs
3196µs
TCA
54035µs
50734µs
47123µs
TM
13µs
TCo
12µs
7µ
TE
4µs
3µs
3µs
TT
54045µs
54043µs
51200µs
Tabela 4.2: Desempenho do Servidor de Áudio
Esse teste mediu os tempos de execução de algumas funções de diferentes versões do
Servidor de Áudio. As versões do Servidor de Áudio utilizadas foram: um Servidor de Áudio
sem codificação que apenas captura o áudio e envia ao cliente que reproduz o áudio recebido;
um Servidor de Áudio com codificação em CPU que captura o áudio, codifica-o utilizando
o codificador LAME original do teste anterior e envia o resultado da codificação ao cliente
que o descodifica e reproduz o áudio resultante; e o Servidor de Áudio com codificação em
GPU, descrito na seção anterior, que utiliza o codificador LAME modificado (com parte do
processamento em GPU).
Os tempos medidos para esse teste foram o tempo de configuração do codificador
LAME (TCL), o tempo de captura do áudio (TCA), o tempo de cópia das amostras para a
memória da GPU (TM), o tempo de codificação (TCo), o tempo de envio dos dados (TE) e o
tempo total de cada ciclo do servidor (TT). Os valores foram medidos em micro segundos (µs).
Com isso obteve-se a Tabela 4.2.
Vale destacar que a captura do áudio está sujeita a variações de acordo com o áudio
introduzido. Logo, o desvio padrão nesse caso chegou a 35% no Servidor de Áudio sem
codificação, a 37,2% no Servidor de Áudio com codificação em CPU e a 62% Servidor de
Áudio com codificação em GPU. Além disso pode-se considerar que o tempo de captura, em
teoria, deve ser semelhante para os diferentes servidores de áudio e, portanto, seus valores
não alteram o resultado desejado, que era medir a eficiência da implementação do codificador
LAME em GPU.
Como demontra a Tabela 4.2, foi possı́vel conseguir uma redução no tempo de envio
dos dados, pois com a codificação a quantidade de dados enviados ao cliente por ciclo9 reduziu
9 Cada
ciclo corresponde ao processo de captura, codificação e envio do áudio.
31
de 2304 bytes10 para aproximadamente 208 bytes. Deve-se lembrar que o tempo de envio
dos dados medido neste teste consiste no tempo de execução da função responsável por essa
ação. Porém, para isso, foi necessário introduzir a configuração do LAME, que precisa ser
reconfigurado para cada codificação.
Da mesma forma, foi possı́vel reduzir o tempo de codificação utilizando a versão modificada do LAME que possui parte de seu processamento em GPU. Nessa comparação, o Servidor de Áudio com codificação em GPU conseguiu codificar o áudio 1,71 vezes mais rápido que
o Servidor de Áudio com codificação em CPU. Porém, mais uma vez, foi necessário acrecentar
o tempo da cópia dos dados do áudio para a memória da GPU. Dessa forma, apesar do desempenho geral do Servidor de Áudio não ter sido melhorado, pois, em proporção ao tempo total
de execução de um ciclo, o tempo de codificação é muito pequeno e somado ao tempo gasto
transferindo os dados para a memória da GPU, ou seja, o ganho de desempenho de codificação
não foi suficiente para compensar o tempo gasto transferindo os dados para a memória da GPU.
Porém, o aumento da quantidade de dados a serem processados podem permitir que o ganho de
desempenho do processamento compense o tempo de cópia dos dados para a memória da GPU.
Mas para comprovar esta afirmação é necessário analisar a variação dos tempos de cópia dos
dados para a memória da GPU e de codificação em relação ao aumento da quantidade de dados.
Ainda existem partes do codificador LAME que podem ser implementadas em CUDA
a fim de melhorar ainda mais seu desempenho. Por exemplo, reescrever as funções que utilizam a FFT11 em CUDA; paralelizar os filtros Replay Gain12 e passa-baixa13 , assim como
outras otimizações na função de análise psicoacústica e nas funções de codificação, entre outras
possı́veis implementações que não foram citadas aqui. Além disso um novo teste medindo o
tempo de execução das funções do LAME que foram implementadas em CUDA seria útil para
analisar o ganho de desempenho obtido pelo codificador.
10 1152
amostras de 2 bytes cada
possı́vel utilizar a biblioteca cuFFT que possui um amplo suporte às funções de FFT.
12 Responsáveis por normalizar o ruı́do perceptı́vel nas amostras de áudio
13 Responsável por atenuar a amplitude das freqüências
11 É
32
5
Conclusão
Neste trabalho foram introduzidos os principais conceitos a respeito da programação
paralela e alguns conceitos que envolvem o áudio digital e o processamento de sinais, com o
objetivo de desenvolver um modelo de Servidor de Áudio que permita o desenvolvimento de
uma aplicação que utilize o processamento em GPU para melhorar o desempenho de técnicas
de processamento de áudio. Demostrando, com isso, que o modelo de programação paralela em
GPU é adequado para ser utilizado em aplicações multimı́dia.
Foi implementado um Servidor de Áudio, baseado no modelo proposto, que utiliza a
programação em GPU para implementar a codificação do áudio. A idéia foi utilizar a codificação
do áudio para reduzir a quantidade de dados enviada e reduzir o tempo de codificação utilizando
uma implementação do codificador LAME em CUDA. Para avaliar o Servidor de Áudio implementado, foram realizados experimentos. O primeiro experimento comparou o desempenho
do codificador de áudio em GPU com o desempenho de sua versão em CPU. Enquanto o segundo experimento comparou os tempos de execução das funções de três versões do Servidor de
Áudio: uma versão sem codificação, uma com codificação em CPU e a última com codificação
em GPU.
Os experimentos demostraram que a aplicação implementada não possuia volume de
dados intenso ou grau de paralelismo suficientes para tornar o resultado do uso da técnica de
processamento paralelo em GPU expressivo. Porém, os resultados mostraram que o modelo de
programação paralela em GPU foi capaz de otimizar o desempenho da codificação do áudio,
como foi proposto.
Assim, foi possı́vel demonstrar que a utilização da GPU para a programação de propósito
33
geral está evoluindo rapidamente de maneira a se tornar um meio eficiente e viável para a
implementação de diversos tipos de aplicações. E que os sistemas multimı́dia podem se aproveitar do aumento da capacidade de processamento, proporcionado pelos modelos de programação
em GPU, para melhorar do desempenho das técnicas existentes e para a criação de novas
técnicas.
34
Referências Bibliográficas
BORN, G. Formats Handbook. London: Thomson Computer Press, 1995.
CAVAGNOLO, B.; BIER, J. Introduction to digital audio compression.
CHONG, J. et al. Data-parallel large vocabulary continuous speech recognition on graphics
processors. In: Proceedings of the 1st Annual Workshop on Emerging Applications and Many
Core Architecture (EAMA). [S.l.: s.n.], 2008. p. 23–35.
DUHAMEL, P.; VETTERLI, M. Fast fourier transforms: A tutorial review and a state of
the art. Signal Process., Elsevier North-Holland, Inc., Amsterdam, The Netherlands, The
Netherlands, v. 19, n. 4, p. 259–299, 1990. ISSN 0165-1684.
HALFHILL, T. R. Parallel Processing with CUDA. January 28 2008. InStat Microprocessor
Report.
HENNESSY, J. L.; PATTERSON, D. A. Computer Architecture; A Quantitative Approach.
San Francisco, CA, USA: Morgan Kauffman Publishers Inc., 1990. ISBN 1558600698.
HUFFMAN, D. A. A method for the construction of minimum-redundancy codes. Proceedings of the IRE, v. 40, n. 9, p. 1098–1101, 1952. Disponı́vel em:
<http://ieeexplore.ieee.org/xpls/abs all.jsp?arnumber=4051119>.
LAGO, N. P. Processamento Distribuı́do de Áudio em Tempo Real. Abril 2004.
MANAVSKI., S. A. Cuda compatible gpu as an efficient hardware accelerator for aes
cryptography. In: . [S.l.: s.n.], 2007. p. 65–68.
NVIDIA. NVIDIA CUDA Compute Unified Device Architecture Programming Guide. Version
2.0. June 7 2008.
PAN, D. Y. Digital audio compression. Digital Tech. J., Digital Equipment Corp., Acton, MA,
USA, v. 5, n. 2, p. 28–40, 1993. ISSN 0898-901X.
PAN, D. Y. A tutorial on mpeg/audio compression. IEEE MultiMedia, IEEE Computer Society
Press, Los Alamitos, CA, USA, v. 2, n. 2, p. 60–74, 1995. ISSN 1070-986X.
PEDDIE, J. Digital Media Technology: Industry Trends and Developments. 2001. IEEE
Computer Graphics and Applications.
PEERCY, M.; SEGAL, M.; GERSTMANN, D. A performance-oriented data parallel virtual
machine for gpus. In: SIGGRAPH ’06: ACM SIGGRAPH 2006 Sketches. New York, NY, USA:
ACM, 2006. p. 184. ISBN 1-59593-364-6.
POHLMANN, K. C. Principles of Digital Audio. [S.l.]: McGraw-Hill Professional, 2000.
ISBN 0071348190.
35
POSTEL, J. User Datagram Protocol. [S.l.], August 1980. 3 p. Disponı́vel em:
<http://www.rfc-editor.org/rfc/rfc768.txt>.
PRINCEN, J. P.; BRADLEY, A. B. Analysis/synthesis filter bank design based on time domain
aliasing cancellation. IEEE Transaction on Acoustics, Speech and Signal Processing, n. 5, p.
1153–1161, out. 1986.
ROADS, C. et al. The Computer Music Tutorial. Cambridge, MA, USA: MIT Press, 1996.
ISBN 0-252-18158-3.
SILBERSCHATZ, A.; GALVIN, P. B. Operating System Concepts. New York, NY, USA: John
Wiley & Sons, Inc., 2000. ISBN 0471418846.
VAIDYANATHAN, P. P. Quadrature mirror filter banks, M-band extensions and perfect
reconstruction techniques. v. 4, n. 3, p. 4–20, jul. 1987.
WATKINSON, J. Introduction to Digital Audio. Newton, MA, USA: Butterworth-Heinemann,
2001. ISBN 0240516435.
36
ANEXO A -- CUDA Application Programming Interface
A interface de programação (API) do modelo CUDA fornece um meio de programadores familiarizados com a linguagem C escreverem facilmente programas para executar em
GPU. Para isso um conjunto mı́nimo de extensões da linguagem C permitem indicar partes do
código para ser executado em GPU. A API consiste também em uma biblioteca de execução1
que permite o controle de mais de uma GPU a partir da CPU; funções especı́ficas para executarem em GPU; e uma componente comum que define tipos e um subconjunto da biblioteca
padrão C permitem a execução tanto em CPU quanto em GPU. As únicas funções da biblioteca
padrão C suportadas para executar em GPU são as disponibilizadas pela componente comum
da biblioteca de execução.
A.1
Extensões da Linguagem C
A API CUDA define 4 extensões para a linguagem C: qualificadores de função, que
definem se a função deve ser chamada em CPU ou GPU e onde ela deve ser executada; qualificadores de variáveis que especificam em qual memória a variável será alocada; uma nova diretiva
que especifica como um kernel deve ser executado; e variáveis pré-definidas que especificam as
dimensões do grid e dos blocos e os ı́ndices dos blocos e threads.
A.1.1
Qualificadores de Função
Os qualificadores de função device , global e host são responsáveis por de-
finir se a função definida será executada em CPU ou GPU e a partir de onde ela pode ser
1 CUDA
runtime library
37
invocada.
O qualificador
device
declara uma função que é executada em GPU e pode ser
chamada apenas a partir da GPU. O qualificador
global
declara uma função como sendo
um kernel. Tal função é executada em GPU e pode ser chamada apenas a partir da CPU. O
qualificador host declara uma função que é executada em CPU e pode ser chamada apenas a
partir da CPU. É equivalente declarar uma função com o qualificador host ou sem nenhum
dos qualificadores de função. O qualificador
host
pode ser utilizado em conjunto com o
qualificador device . Nesse caso o código será compilado para GPU e CPU.
Os qualificadores de função possuem restrições. Funções definidas para executar em
GPU não suportam recursão; não podem possuir declarações de variáveis estáticas; e não podem ter número variável de argumentos. Ponteiros para funções device não são suportados.
Os qualificadores
global
e
host
não podem ser usados juntos. Funções
global
de-
vem retornar void, pois essas funções são assı́ncronas, ou seja, a chamada retorna antes do fim
de sua execução. Parâmetros de funções
global
são passados para a GPU pela memória
compartilhada e são limitados a 256 bytes.
A.1.2
Qualificadores de Variáveis
Os qualificadores de variáveis definem em qual tipo de memória a variável declarada
será alocada.
Varı́aveis declaradas com o qualificador
device
serão alocadas na memória global
da GPU; que permanecem em memória durante o tempo de execução do kernel; e acessı́veis por
todas as threads do grid e a partir da CPU através da biblioteca de execução. Qualquer um dos
outros qualificadores de variáveis podem ser utilizados junto com o qualificador device . O
qualificador constant define variáveis que residem na memória constante; que permanecem
em memória durante o tempo de execução do kernel; e acessı́veis por todas as threads do grid
e a partir da CPU através da biblioteca de execução. O qualificador shared define variáveis
alocadas na memória compartilhada de um bloco de threads; que permanecem em memória
38
durante a execução do bloco; e acessı́veis a todas as threads do bloco. Apenas após a execução
do comando
syncthreads() que a escrita a variáveis compartilhadas são garantidas de serem
vistas pelas outras threads do bloco.
Os qualificadores de variáveis possuem restrições. Não é permitido utilizá-los em
struct ou union, em parâmetros formais e em variáveis locais de uma função que executa
em CPU. Os qualificadores
shared
e
constant
implicam em armazenamento estático.
Variáveis device e constant são declaradas fora de funções. Variáveis constant não
podem ser definidas em GPU, apenas através de funções de execução especı́ficas em CPU.
Variáveis
shared
não podem possuir uma declaração como parte de suas declarações. Os
endereços obtidos de variáveis device , shared ou constant podem ser utilizadas apenas em GPU. Entretanto, os endereços de variáveis
device
ou
constant
obtidos através
da função cudaGetSymbolAddress()2 podem ser usadas em CPU.
A.1.3
Parâmetros de Configuração da Execução
Qualquer chamada de uma função
global
deve especificar uma configuração de
execução para a chamada.
A configuração de execução define a dimensão do grid e dos blocos que serão usados
para executar a função no dispositivo. A especificação da configuração é feita inserindo uma
expressão da forma <<< Dg, Db, Ns, S >>> entre o nome da função e a lista de argumentos.
Dg é do tipo dim3, que será abordado na seção A.2.1, e especifica o tamanho do grid em
até duas dimensões, onde o número de blocos a serem lançados é Dg.x * Dg.y (o número de
blocos na dimensão x vezes o número de blocos na dimensão y). Db também é do tipo dim3
e especifica o tamanho de cada bloco em até três dimensões, tal que Db.x * Db.y * Db.z (o
número de blocos na dimensão x vezes o número de blocos na dimensão y vezes o número de
blocos na dimensão z) equivale ao número de threads por bloco. Ns é do tipo size t e especifica
o número de bytes na memória compartilhada que será alocada dinamicamente em cada bloco
para uma chamada além da memória alocada estaticamente. Esse valor é usado para definir o
2 Seção
4.5.2.3 do Guia de Programação CUDA [NVIDIA, 2008]
39
tamanho de vetores alocados dinamicamente. O argumento Ns é opcional e tem valor padrão
igual a 0. S é do tipo cudaStream t e especifica o stream associado. O argumento S é opcional
e tem valor padrão igual a 0.
Os argumentos de configuração de execução são avaliados antes dos argumentos da
função e ambos são passados através da memória compartilhada para a GPU. Se algum dos
parâmetros de configuração da execução forem maior que o permitido a execução irá falhar.
A.1.4
Variáveis Pré-definidas
As variáveis pré-definidas são variáveis definidas automaticamente a partir da chamada
de uma função. Elas possuem as dimensões e tamanhos do grid e dos blocos e os ı́ndices dos
blocos e threads.
A variável gridDim é do tipo dim3 e contém o tamanho do grid para todas as suas
dimensões. A variável blockIdx é do tipo uint3, que será explicado na seção A.2.1, e contém o
ı́ndice do bloco corrente para cada uma das dimensões do grid. A variável blockDim é do tipo
dim3 e contém o tamanho do bloco todas as suas dimensões. A variável threadIdx é do tipo
uint3 e contém o ı́ndice da thread corrente para cada uma das dimensões do bloco. E a variável
warpSize é do tipo int e contém o tamanho do warp em threads.
Não é possı́vel indicar o endereço ou atribuir valor a nenhuma dessas variáveis.
A.1.5
O Compilador NVCC
O compilador nvcc busca simplificar o processo de compilação do código CUDA. O
compilador provê opções de linha de comando simples e familiares.
A rotina básica do nvcc consiste em separar o código GPU do código CPU e compilar
o código GPU em uma forma binária conhecida como cubin. O código CPU gerado permanece
em C e será compilado com outra ferramenta no último estágio da compilação.
Aplicações podem ignorar o código CPU gerado e carregar e executar o código cubin
40
em GPU diretamente utilizando a API do driver3 ou podem linkar o código cubin com o código
CPU.
O código CUDA é compilado de acordo com as regras de sintaxe da linguagem C++.
C++ é totalmente suportado no código CPU, no entanto, apenas o subconjunto de regras de
C são totalmente suportadas em GPU. Caracterı́sticas especı́ficas como classes, herança, ou
declaração de variáveis em blocos básicos não são suportadas. Como conseqüência do uso das
regras de sintaxe de C++, ponteiros void não podem ser associados a ponteiros não-void sem o
uso de typecast.
O nvcc introduz duas diretivas:
noinline e # pragma unroll.
Por padrão, uma funçao device é definida como inline, ou seja, a função é copiada
inteira para cada posição onde ela é chamada. A diretiva
noinline
é utilizada para indicar
para o processador não fazer isso se possı́vel. Ainda assim, a função deve estar no mesmo
arquivo em que ela é chamada.
Por padrão o compilador desenrola pequenos loops para melhorar o desempenho da
aplicação. A diretiva # pragma unroll permite controlar o desenrolamento de qualquer loop.
A diretiva deve ser inserida imediatamente antes do loop. Ela pode ser opcionalmente seguida
pelo número de vezes que o loop será desenrolado. # pragma unroll 1 indica ao compilador
que o loop não deve ser desenrolado. Se o número de vezes que o loop será desenrolado, o
compilador desenrola o loop todo.
A.2
Componente de Execução Comum
A componente de execução comum4 , como o nome diz, pode ser usada tanto em GPU
quanto em CPU.
3 Ver
seção A.4.2
Runtime Component
4 Common
41
A.2.1
Tipos Pré-definidos
char1, uchar1, char2, uchar2, char3, uchar3, char4, uchar4, short1, ushort1,
short2, ushort2, short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4,
uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, float1, ufloat1, float2,
ufloat2, float3, ufloat3, float4, ufloat4 são estruturas baseadas dos tipos básicos inteiro e ponto
flutuante. Suas primeira, segunda, terceira e quarta componentes são acessı́veis através dos
campos x, y, z e w respectivamente. Todos esses tipos possuem um construtor com a forma
make <nome do tipo>(). Por exemplo: make int2(int x, int y) cria uma variável do tipo int2
com valor (x, y).
O tipo dim3 é baseado no tipo uint3. Esse tipo é usado para especificar dimensões.
Quando uma variável é definida com esse tipo, qualquer componente não especificada possuirá,
por padrão, valor 1.
A.2.2
Funções Matemáticas
As funções matemáticas da biblioteca padrão C/C++ suportadas em CUDA são espe-
cificadas na componente de execução comum, ou seja, podem ser executadas tanto em CPU
como em GPU.
A Seção B.1 do Guia de Programação CUDA [NVIDIA, 2008] contém uma lista das
funções matemáticas da biblioteca padrão C/C++ que são suportadas em CUDA.
A.2.3
Funções de Tempo
Quando a função clock t clock(); é executada em GPU, ela retorna o valor de um
contador do multiprocessador que é incrementado a cada ciclo do clock. Cada multiprocessador
da GPU possui um contador individual. Coletando o valor desse contador no inı́cio e no final
de um kernel, tirando a diferença entre as duas coletas e guardando o resultado por thread é
possı́vel medir o número de clocks que cada thread precisou para completar a execução do
42
kernel.
A.2.4
Tipo Textura
A biblioteca de execução CUDA suporta um subconjunto de instruções para o acesso
a memória de textura. É possı́vel obter benefı́cios de desempenho ao ler dados da memória de
textura ao invés da memória global.
A memória de textura é lida através de um kernel usando funções chamadas de fetches de textura. O primeiro parâmetro de um fetch é um objeto chamado de referência de
textura. Uma referência de textura define qual parte da textura será buscada. Ela deve ser limitada através da CPU a algumas regiões de memória, chamadas texturas, antes de serem usadas
pelo kernel. Várias regiões distintas podem ser limitadas a uma mesma textura ou a texturas
sobrepostas na memória.
Uma referência de textura possui vários atributos. Um deles é a dimensão da textura,
que especifica onde a textura está endereçada. Uma textura é endereçada como um vetor de até
3 dimensões. Os elementos do vetor são chamados de elementos de textura.
Outros atributos definem os tipos de entrada e saı́da do fetch de textura; como as coordenadas das entradas são interpretadas e que processamento deve ser feito.
Alguns atributos de textura são imutáveis e devem ser conhecidos em tempo de compilação.
Eles são especificados ao declarar a referência de textura. Uma referência de textura é declarada
no escopo do arquivo como uma variável do tipo textura: texture<Type, Dim, ReadMode>
texRef;. Onde Type especifica o tipo do dado que será retornado pela busca da textura. Type é
restrito a tipos inteiros básicos, ponto flutuantes de precisão simples e qualquer dos tipos com
1, 2 ou 4 componentes definidos na seção A.2.1. Dim especifica a quantidade de dimensões
da referência de textura e pode possuir valor igual a 1, 2 ou 3. Dim é um argumento opcional e, por padrão, possui valor igual a 1. ReadMode possui valor igual a cudaReadModeElementType ou cudaReadModeNormalizedFloat. Caso ReadMode seja cudaReadModeNormalizedFloat e Type seja um inteiro de 16 ou 8 bits, o valor retornado é convertido para
43
ponto flutuante. Caso ReadMode seja cudaReadModeElementType, nenhuma conversão é
realizada. O atributo ReadMode é opcional e possui como valor padrão cudaReadModeElementType.
Os outros atributos de uma referência de textura são mutáveis e podem ser alteradas
em tempo de execução através de instruções da CPU. Eles especificam onde as coordenadas da
textura estão normalizadas ou não; o modo de endereçamento; e filtros de textura.
A.3
Componente de Execução em GPU
As funções da componente de execução em GPU podem ser utilizadas apenas em
funções que executam na GPU.
A.3.1
Funções Matemáticas
Para algumas das funções referenciadas na seção A.2.2 existe uma versão menos pre-
cisa porém mais rápida na componente de execução em GPU. Seus nomes possuem o mesmo
prefixo com
(por exemplo:
sinf(x)). Essas funções estão listadas na Seção B.2 do Guia de
Programação CUDA [NVIDIA, 2008].
O compilador também possui uma opção (-use fast math) para forçar cada função a
compilar a sua versão menos precisa, se ela existir.
A.3.2
Função de Sincronização
A função
syncthreads() sincroniza todas as threads em um bloco. Uma vez que
todas as threads atingiram este ponto, a execução prossegue normalmente.
É permitido utilizar a função
syncthreads() instruções condicionais, apenas se as
condições forem avaliadas identicamente para todas as threads do bloco. Se essa condição não
for atendida, a execução pode resultar efeitos não desejados.
44
A.3.3
Funções de Textura
Uma textura pode ser qualquer região de memória linear ou um vetor CUDA (que são
regiões de memória otimizadas para utilizar texturas).
Utilizando uma região de memória linear, a textura pode ser acessı́vel através da
famı́lia de funções tex1Dfetch().
Quando vetores CUDA são usados para acessar uma textura, utiliza-se as funções
tex1D(), tex2D ou tex3D.
A.3.4
Funções Atômicas
Uma função atômica realiza leitura, modificação e escrita em uma única operação em
uma palavra de 32 ou 64 bits em algum endereço na memória global ou na compartilhada. Por
exemplo, a função atomicAdd() lê uma palavra de 32 bits de algum endereço de memória na
memória global na compartilhada, soma um inteiro à palavra e escreve o resultado no mesmo
endereço. A operação é atômica no sentido de garantir que será executada sem interferência de
outras threads.
A.4
Componente de Execução em CPU
A componente de execução em CPU fornece um conjunto de funções para manipular o
gerenciamento da GPU; o gerenciamento de contexto; o gerenciamento de memória; o controle
de execução; o gerenciamento das referências de textura; e a interoperabilidade com OpenGL
e Direct3D. A componente de execução em CPU é composta de duas partes: a API do driver
CUDA e a API de execução CUDA. Essas partes são mutuamente exclusivas, ou seja, só é
permitido a uma aplicação usar uma delas.
45
A.4.1
API de Execução
Não há uma inicialização explı́cita para a API de execução. Ela é inicializada ao exe-
cutar a primeira função da biblioteca de execução.
Para gerenciar as GPUs instaladas no sistema, são utilizadas as funções cudaGetDeviceCount(), cudaGetDeviceProperties() e cudaSetDevice(). cudaGetDeviceCount() e cudaGetDeviceProperties() fornecem um meio de enumerar as GPUs e obter suas informações.
Enquanto cudaSetDevice é usada para definir a GPU que será associada a thread da CPU.
Regiões de memória linear são alocadas usando a função cudaMalloc() ou cudaMallocPitch() e liberadas usando cudaFree(). Vetores CUDA são alocados pela função cudaMallocArray() e liberadas pela função cudaFreeArray(). cudaMallocArray() exige que uma
descrição de formato seja criada através da função cudaCreateChannelDesc(). Um endereço
de variável alocada ne memória global pode ser obtido pela função cudaGetSymbolAddress().
E o tamanho da memória alocada é obtido pela função cudaGetSymbolSize().
O Gerenciamento da referência de textura utiliza o tipo textureReference para definir
uma referência de textura. Antes que um kernel possa usar uma referência de textura para ler da
memória de textura, a referência de textura deve ser vinculada a uma textura usando a função
cudaBindTexture() ou a cudaBindTextureToArray(). E para desvincular uma referência de
textura, utiliza-se a função cudaUnbindTexture().
Para a interoperabilidade com o OpenGL, um buffer deve ser registrado em CUDA
antes que ele possa ser mapeado. Isso é feito com a função cudaGLRegisterBufferObject().
Uma vez registrado, o buffer pode sr lido ou escrito por kernels usando o endereço de memória
da GPU retornado por cudaGLMapBufferObject(). E para eliminar o mapeamento e o registro
utiliza-se, respectivemente, as funções cudaGLUnmapBufferObject() e cudaGLUnregisterBufferObject().
E para a interoperabilidade com o Direct3D é necessário determinar qual GPU será utilizada antes de qualquer execução. Isso é feito através da função cudaD3D9SetDirect3DDevice().
46
Os recursos Direct3D são registrados em CUDA pela função cudaD3D9RegisterResources().
Para remover o registro utiliza-se cudaD3DUnregisterVertexBuffer(). Assim que os recursos
foram registrados, eles podem ser mapeados em CUDA quantas vezes forem necessárias através
da função cudaD3D9MapResources(). E da mesma forma, para eliminar o mapeamento
utiliza-se cudaD3D9UnmapResources(). Um recurso mapeado pode ser lido ou escrito por
kernels utilizando o endereço de memória retornado por cudaD3D9ResourceGetMappedPointer()
e o tamanho e o pitch são retornados pelas funções: cudaD3D9ResourceGetMappedSize(),
cudaD3D9ResourceGetMappedPitch() e cudaD3D9ResourceGetMappedPitchSlice().
A.4.2
API do Driver
Antes de executar qualquer função da API do driver, é necessário uma inicialização
com a função cuInit().
Para manipular as GPUs instaladas no sistema, são utilizadas, entre outras, as funções
cuDeviceGetCount() e cuDeviceGet().
A API do driver também permite a manipulação do contexto. Todos os recursos e
ações realizados através da API do driver são encapsuladas em um contexto CUDA. Uma thread de CPU pode possuir apenas um contexto ativo por vez. Para criar um contexto, usa-se
a função cuCtxCreate(). Cada thread da CPU possui uma pilha de contextos. cuCtxCreate()
põe um novo contexto no topo da pilha. Para desvincular o contexto ativo de uma thread da
CPU, a função cuCtxPopCurrent(). Se houver algum contexto anterior, ele será reativado.
Uma contagem de uso é mantida para cada contexto. A função cuCtxCreate() cria um contexto com contagem igual a 1 que é incrementado pela função cuCtxAttach() e decrementado
pela função cuCtxDetach(). Um contexto é destruı́do quando a contagem chega a 0 ao chamar
a função cuCtxDetach() ou cuCtxDestroy(). Clientes da API podem usar as funções cuCtxPushCurrent() e cuCtxPopCurrent() para criar contextos.
O controle de execução pode ser manipulado pelas seguintes funções. A função cuFuncSetBlockShape() que define o número de threads por bloco para uma dada função. cu-
47
FuncSetSharedSize() define o tamanho de memória compartilhada para a função. A famı́lia
de funções cuParam*() é usada para especificar os parâmetros do kernel que será invocado.
cuLaunchGrid() ou cuLaunch() são utilizadas para invocar um kernel.
O gerenciamento de memória é feito através das funções cuMemAlloc(), cuMemAllocPitch() e cuMemFree() para manipular memória linear. E através das funções cuArrayCreate() e cuArrayDestroy() para manipular vetores CUDA.
Para utilizar a memória de textura, uma referência de textura deve ser criada utilizando
a função cuTexRefSetAddress() ou cuTeXRefSetArray().
A interoperabilidade com o OpenGL requer uma inicialização pela função cuGLInit().
Após isso, um buffer deve ser registrado utilizando a função cuGLRegisterBufferObject()
e mapeado com a função cuGLMapBufferObject(). Para eliminar o mapeamento, utilizase a função cuGLUnmapBufferObject(). E a função cuGLUnregisterBufferObject() para
remover o registro o buffer.
E para a interoperabilidade com o Direct3D é necessária a criação de um contexto. Isso
pode ser feito pela função cuD3D9CtxCreate(). Os recursos Direct3D podem ser registrados
em CUDA usando a função cuD3D9RegisterResource(). Esse registro pode ser eliminado pela
função cuD3D9UnregisterVertexBuffer(). Após os recursos serem registrados em CUDA, eles
podem ser mapeados pela função cuD3D9MapResources(). E o mapeamento pode ser eliminado pela função cuD3D9UnmapResources(). Um recurso mapeado pode ser lido e escrito
pelo kernel através do ponteiro retornado pela função cuD3D9ResourceGetMappedpointer(),
do tamanho retornado pela função cuD3D9ResourceGetMappedSize() e do pitch retornado
pelas funções cuD3D9ResourceGetMappedPitch() e cuD3D9ResourceGetMappedPitchSlice().
48
ANEXO B -- Speaker (servidor)
/*******************************************************************************
* Este codigo foi escrito por mim e tem a finalidade de capturar o audio
* codifica-lo de forma paralela e envia-lo a um cliente
******************************************************************************/
/* Bibliotecas Gerais */
#include <cuda.h>
#include <cuda_runtime.h>
#include <lhal04.h>
#include <lhal04_lame/include/lame.h>
void reusePort(int s){
int one=1;
if ( setsockopt(s,SOL_SOCKET,SO_REUSEADDR,(char *) &one,sizeof(one)) == -1 ){
printf("error in setsockopt,SO_REUSEPORT \n");
exit(-1);
}
}
49
int get_gpu_buffer_size( lame_global_flags *gf, int bytesPerSample, int pad )
{
int nsamp = lame_get_num_samples(gf);
return
nsamp*bytesPerSample + pad;
//assume 2 bytes per sample (16-bit)
}
int main(int argc, char **argv){
/* Variaveis para manipulacao do socket */
int sd;
struct sockaddr_in server;
struct hostent *hp, *gethostbyname();
struct servent *sp;
struct sockaddr_in to;
int to_len;
int length;
char localhost[MAXHOSTNAME];
char msg[MAXHOSTNAME];
/* Variaveis para manipulacao do dispositivo de audio */
int rc,i;
snd_pcm_t *handle = NULL;
snd_pcm_hw_params_t *params;
unsigned int val;
50
int direction;
snd_pcm_uframes_t frames;
buffer_t *buffer;
int size;
int num_samples_read;
lame_global_flags *gfp;
unsigned char *mp3buffer;
float *gpu_buffers[2];
int gpu_buffer_sz;
float in_buffer_l[1152];
float in_buffer_r[1152];
struct timeval tt1,tt2;
struct timeval tcl1,tcl2;
struct timeval tca1,tca2;
struct timeval ta1,ta2;
struct timeval tco1,tco2;
struct timeval te1,te2;
/* Configuracao do Socket */
if(argc > 2){
fprintf(stderr, "Uso correto: speaker <porta>\n");
exit(1);
}
sp = getservbyname("echo", "udp");
51
/* get Host information, NAME and INET ADDRESS */
gethostname(localhost, MAXHOSTNAME);
printf("----Speaker running at host NAME: %s\n", localhost);
if
( (hp = gethostbyname(localhost)) == NULL ) {
fprintf(stderr, "Can’t find host name\n");
exit(-1);
}
bcopy ( hp->h_addr, &(server.sin_addr), hp->h_length);
printf("
(Speaker INET ADDRESS is: %s )\n", inet_ntoa(server.sin_addr));
/* Construct name of socket to send to. */
server.sin_family = AF_INET;
server.sin_addr.s_addr = htonl(INADDR_ANY);
if (argc == 1)
server.sin_port = htons(0);
else
server.sin_port = htons(atoi(argv[1]));
/* Create socket on which to send
and receive */
sd = socket (AF_INET,SOCK_DGRAM,0);
/* to allow another process to use the same port
howver, only ONE gets the message */
reusePort(sd);
52
if ( bind( sd, (struct sockaddr *) &server, sizeof(server) ) < 0 ) {
close(sd);
perror("binding name to datagram socket");
exit(-1);
}
/* get port information and
prints it out */
length = sizeof(server);
if ( getsockname (sd, (struct sockaddr *)&server,&length) ) {
perror("getting socket name");
exit(0);
}
printf("Server Port is: %d\n", ntohs(server.sin_port));
to_len = sizeof(to);
printf("\n...server is waiting...\n");
if ((rc=recvfrom(sd, msg, sizeof(msg), 0, (struct sockaddr *) &to, &to_len)) <
0)
perror("receiving datagram
message");
printf("Enviando audio para: %s:%d\n", inet_ntoa(to.sin_addr),
htons(to.sin_port));
if ((hp = gethostbyaddr((char *)&to.sin_addr.s_addr,
sizeof(to.sin_addr.s_addr), AF_INET)) == NULL)
fprintf(stderr, "Can’t find host %s\n", inet_ntoa(to.sin_addr));
53
/* Configuracao do dispositivo de audio */
/* Abre o dispositivo PCM "default" para captura (gravacao) */
rc = snd_pcm_open(&handle, "default", SND_PCM_STREAM_CAPTURE, 0);
if (rc < 0){
fprintf(stderr, "Nao consegui abrir o dispositivo pcm: %s\n",
snd_strerror(rc));
exit(1);
}
/* Aloca um objeto de parametros do harware */
snd_pcm_hw_params_malloc(&params);
/* Preenche os parametros com valores default */
snd_pcm_hw_params_any(handle, params);
/* Define os parametros de hardware desejados */
printf("Configuracoes de Audio...\n");
/* Define modo de acesso pcm como sendo entrelacado (interleaved) */
snd_pcm_hw_params_set_access(handle, params, SND_PCM_ACCESS_RW_INTERLEAVED);
/* Define o formato do audio como signed 16 bits little-endian */
snd_pcm_hw_params_set_format(handle, params, SND_PCM_FORMAT_S16_LE);
/* Define dois canais (stereo) */
snd_pcm_hw_params_set_channels(handle, params, CANAIS);
54
/* Define a taxa de amostragem aproximada como 44100 bits/s (qualid. de CD) */
val = SAMPLE_RATE;
direction = 0;
snd_pcm_hw_params_set_rate_near(handle, params, &val, &direction);
/* Define o tamanho do periodo em frames */
frames = 32;
direction = 0;
snd_pcm_hw_params_set_period_size_near(handle, params, &frames, &direction);
/* Escreve os parametros no Driver */
rc = snd_pcm_hw_params(handle, params);
if (rc < 0){
fprintf(stderr, "\nNao consegui definir os parametros de hw: %s\n",
snd_strerror(rc));
exit(1);
}
snd_pcm_hw_params_get_channels(params, &val);
printf("Numero de canais: %d\n",val);
snd_pcm_hw_params_get_rate(params, &val, &direction);
printf("Taxa de Amostragem: %d Hz\n",val);
printf("...OK\n\n");
gfp = lame_init();
lame_set_num_channels(gfp,CANAIS);
lame_set_mode(gfp,MONO);
55
lame_set_in_samplerate(gfp,SAMPLE_RATE);
lame_set_num_samples(gfp,1152);
lame_init_params(gfp);
lame_print_config(gfp);
lame_print_internals(gfp);
size = lame_get_size_mp3buffer(gfp);
buffer = (buffer_t *) malloc(NSAMPLES*2);
mp3buffer = (unsigned char *) malloc(size);
printf("Tamanho do buffer pcm: %d amostras\n",NSAMPLES);
printf("Tamanho do buffer mp3: %d bytes\n\n",size);
lame_close(gfp);
/* Loop principal */
while(1){
gfp = lame_init();
lame_set_num_channels(gfp,CANAIS);
lame_set_mode(gfp,MONO);
lame_set_num_samples(gfp,NSAMPLES);
lame_set_in_samplerate(gfp,SAMPLE_RATE);
lame_set_num_samples(gfp,1152);
rc = lame_init_params(gfp);
/* Captura Audio */
56
rc = snd_pcm_readi(handle, buffer, NSAMPLES*2);
if (rc == -EPIPE){
fprintf(stderr, "Overrun ocurred\n");
snd_pcm_prepare(handle);
rc = NSAMPLES*2;
}
else if (rc < 0){
fprintf(stderr, "Erro na leitura: %s\n", snd_strerror(rc));
}
else if (rc != NSAMPLES*2){
fprintf(stderr, "Short read. Expected %d samples, read %d samples\n",
NSAMPLES, rc/2);
/* As amostras sao de 2 bytes cada, por isso eh necessario dividir o
* tamanho do buffer e o rc por 2, uma vez que estas variaveis sao medidas
* em bytes, nao em amostras. */
}
num_samples_read = rc/2;
gpu_buffer_sz = get_gpu_buffer_size( gfp, sizeof(float), 0);
cudaMalloc((void *)&(gpu_buffers[1]), gpu_buffer_sz);
cudaMalloc((void *)&(gpu_buffers[0]), gpu_buffer_sz);
for ( i = 0; i < num_samples_read; i++){
// valores esperados: +/- 32768.0
in_buffer_l[i] = (float) buffer[i];
57
in_buffer_r[i] = (float) buffer[i];
}
cudaMemcpy( &gpu_buffers[0], in_buffer_l,
cudaMemcpy( &gpu_buffers[1], in_buffer_r,
rc = lame_encode_buffer(gfp, &gpu_buffers[0], &gpu_buffers[1], num_samples_read
rc = lame_encode_flush(gfp, mp3buffer, sizeof(mp3buffer));
sendto(sd, mp3buffer, size, 0, (struct sockaddr *) &to, sizeof(to));
lame_close(gfp);
}
}
58
ANEXO C -- Listener (cliente)
/*******************************************************************************
* Este codigo foi escrito por mim e tem a finalidade de receber o audio de um
* servidor, descodifica-lo e reproduzi-lo
******************************************************************************/
/* Bibliotecas Gerais */
#include <lhal04.h>
#include <orig/include/lame.h>
int main(int argc, char **argv){
/* Variaveis para manipulacao do socket */
int sd;
struct sockaddr_in server;
struct hostent *hp, *gethostbyname();
struct servent *sp;
struct sockaddr_in from;
struct sockaddr_in addr;
int fromlen;
int cc;
59
char localhost[MAXHOSTNAME];
char msg[MAXHOSTNAME];
/* Variaveis para manipulacao do audio */
int rc;
snd_pcm_t *handle;
snd_pcm_hw_params_t *params;
unsigned int val;
int direction;
snd_pcm_uframes_t frames;
buffer_t *buffer;
lame_global_flags *gfp;
unsigned char *mp3buffer;
/* Configuracao do Socket */
if(argc != 3){
fprintf(stderr, "Uso correto: listener <end_servidor> <porta>\n");
exit(1);
}
sp = getservbyname("echo", "udp");
/* get Listener Host information, NAME and INET ADDRESS */
gethostname(localhost, MAXHOSTNAME);
60
printf("----Listener running at host NAME: %s\n", localhost);
if
( (hp = gethostbyname(localhost)) == NULL ) {
fprintf(stderr, "Can’t find host %s\n", argv[1]);
exit(-1);
}
bcopy ( hp->h_addr, &(server.sin_addr), hp->h_length);
printf("(Listener INET ADDRESS is: %s )\n", inet_ntoa(server.sin_addr));
/* get Speaker Host information, NAME and INET ADDRESS */
if
( (hp = gethostbyname(argv[1])) == NULL ) {
addr.sin_addr.s_addr = inet_addr(argv[1]);
if ((hp = gethostbyaddr((char *)&addr.sin_addr.s_addr,
sizeof(addr.sin_addr.s_addr),AF_INET)) == NULL) {
fprintf(stderr, "Can’t find host %s\n", argv[1]);
exit(-1);
}
}
printf("----Speaker running at host NAME: %s\n", hp->h_name);
bcopy ( hp->h_addr, &(server.sin_addr), hp->h_length);
printf("(Speaker INET ADDRESS is: %s )\n", inet_ntoa(server.sin_addr));
/* Construct name of socket to send to. */
server.sin_family = AF_INET;
server.sin_port = htons(atoi(argv[2]));
/* Create socket on which to send
and receive */
61
sd = socket (hp->h_addrtype,SOCK_DGRAM,0);
if (sd<0) {
perror("opening datagram socket");
exit(-1);
}
/* Comunica com o Servidor */
strcpy(msg, hp->h_name);
if (sendto(sd, msg, strlen(msg), 0, (struct sockaddr *)&server, sizeof(server)) <
perror("N~
ao consegui comunicar com o servidor");
/* Configuracao do dispositivo de audio */
printf("Configuracoes de Audio...\n");
/*snd_pcm_open(pcm handle, handler identifier, direction, mode)*/
/*mode pode ser bloqueante (0) ou nao bloqueante (1) */
rc = snd_pcm_open(&handle, "default", SND_PCM_STREAM_PLAYBACK, 0);
if (rc < 0){
fprintf(stderr, "Nao consegui abrir o dispositivo pcm: %s\n",
snd_strerror(rc));
exit(1);
}
/* Define os parametros de hardware desejados */
62
snd_pcm_hw_params_malloc(&params);
snd_pcm_hw_params_any(handle, params);
snd_pcm_hw_params_set_access(handle, params, SND_PCM_ACCESS_RW_INTERLEAVED);
snd_pcm_hw_params_set_format(handle, params, SND_PCM_FORMAT_S16_LE);
snd_pcm_hw_params_set_channels(handle, params, CANAIS);
val = SAMPLE_RATE;
snd_pcm_hw_params_set_rate_near(handle, params, &val, &direction);
frames = 32;
snd_pcm_hw_params_set_period_size_min(handle, params, &frames, &direction);
rc = snd_pcm_hw_params(handle, params);
if (rc < 0){
fprintf(stderr, "Nao consegui definir os parametros de hw: %s\n", snd_strerror(
exit(1);
}
snd_pcm_hw_params_get_channels(params, &val);
printf("Numero de canais: %d\n", val);
snd_pcm_hw_params_get_rate(params, &val, &direction);
printf("Taxa de Amostragem: %d Hz\n", val);
printf("...OK\n\n");
gfp = lame_init();
lame_set_num_channels(gfp,CANAIS);
lame_set_mode(gfp,MONO);
63
lame_set_out_samplerate(gfp,SAMPLE_RATE);
lame_set_num_samples(gfp,1152);
lame_set_brate(gfp, 16);
lame_init_params(gfp);
lame_decode_init();
lame_set_decode_only(gfp,1);
lame_print_config(gfp);
lame_print_internals(gfp);
buffer = (buffer_t *) malloc(NSAMPLES*2);
mp3buffer = (unsigned char *) malloc(MAXMP3);
printf("Tamanho do buffer pcm: %d amostras\n",NSAMPLES);
printf("Tamanho do buffer mp3: %d bytes\n\n",MAXMP3);
lame_close(gfp);
/* Loop principal */
while(1){
gfp = lame_init();
lame_set_num_channels(gfp,CANAIS);
lame_set_num_samples(gfp,NSAMPLES);
lame_set_mode(gfp,MONO);
lame_set_num_samples(gfp,1152);
lame_set_brate(gfp, 16);
lame_set_decode_only(gfp,1);
64
rc = lame_init_params(gfp);
lame_decode_init();
/* cc = numero de bytes recebidos */
fromlen = sizeof(from);
cc = recvfrom(sd, mp3buffer, MAXMP3, 0, (struct sockaddr *) &from,
&fromlen);
rc = lame_decode(mp3buffer, MAXMP3, buffer, NULL);
rc = snd_pcm_writei(handle, buffer, NSAMPLES*2);
if (rc < 0)
rc = snd_pcm_recover(handle, rc, 0);
if (rc < 0){
fprintf(stderr, "Reproducao de Audio falhou: %s\n", snd_strerror(rc));
break;
}
else if (rc != NSAMPLES*2)
fprintf(stderr, "Short write. Expected %d samples, written %d samples\n",
NSAMPLES, rc/2);
lame_close(gfp);
}
}
65
ANEXO D -- lhal04.h
/*******************************************************************************
* Biblioteca de Configuracao do Servidor de Audio escrita por mim que é
* incluida no código do servidor e do cliente
******************************************************************************/
/* Bibliotecas Gerais */
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <math.h>
/* Bibliotecas de Audio */
#include <alsa/asoundlib.h>
/* Bibliotecas de Socket */
#include <sys/types.h>
#include <sys/socket.h>
#include <netinet/in.h>
#include <netdb.h>
#include <arpa/inet.h>
#include <sys/time.h>
66
#include <time.h>
#define MAXHOSTNAME 80
/* Usa a API mais recente */
#define ALSA_PCM_NEW_HW_PARAMS_API
typedef short int buffer_t;
/* Define configuracoes de Audio */
#define CANAIS 1
#define SAMPLE_RATE 44100
#define NSAMPLES 1152
#define MAXMP3 lame_get_size_mp3buffer(gfp)
67
ANEXO E -- psyKernel
/*******************************************************************************
* Este codigo é parte das modificaç~
oes feitas no LAME pela NVidia como ponto
* inicial para o concurso de implementaç~
ao do LAME em CUDA. Ele possui os
* Kernels em CUDA e as funcoes em C utilizadas pelo LAME para executar em GPU
******************************************************************************/
#include <stdio.h>
#include "cuda_runtime.h"
#include "cutil.h"
#include "cufft.h"
#include "../include/lame.h"
#include "lame_global_flags.h"
#ifdef __cplusplus
extern "C" void hpf( float *firc , float *samples, float *out, int nCoeff, int
nSamples) ;
extern "C" void scaler( float *data, int num_d, float scale) ;
extern "C" void gpu_init(void);
extern "C" void scaler2( float *d, float *d2, int num_d, float scl);
#else
68
void hpf( float *firc , float *samples, float *out, int nCoeff, int nSamples);
void scaler( float *data, int num_d, float scale);
void gpu_init(void);
void scaler2( float *d, float *d2, int num_d, float scl);
#endif
/* Filtro passa-alta do modelo psicoacústico */
// filter coefficients taken from libmp3lame/pysmodel.c
__constant__ float fircoef[] = {
-8.65163e-18*2, -0.00851586*2, -6.74764e-18*2, 0.0209036*2,
-3.36639e-17*2, -0.0438162 *2, -1.54175e-17*2, 0.0931738*2,
-5.52212e-17*2, -0.313819
*2
};
__global__ void HPFilter(float *firc,
float *firbuf,
float *ns_hpfsmpl,
int szCoeff,
int nSamps ){
int idx = blockIdx.x * gridDim.x + threadIdx.x;
int j = 0;
float sum1;
float sum2;
sum1 = firbuf[idx+10];
sum2 = 0.0f;
69
for( j=0 ; j<10; j+=2 ) {
sum1 += fircoef[j
] * ( firbuf[idx+j
] + firbuf[idx+21-j] );
sum2 += fircoef[j+1] * ( firbuf[idx+j+1 ] + firbuf[idx+21-j-1] );
}
ns_hpfsmpl[idx] = sum1+sum2;
}
static float *gpuBuf = NULL;
static float *hpfBuf = NULL;
void gpu_init(void) {
CUDA_SAFE_CALL(
cudaMalloc((void**)&gpuBuf, 50000)
);
CUDA_SAFE_CALL(
cudaMalloc((void**)&hpfBuf, 50000)
);
}
void hpf(float *coeff ,
float *samples,
float *out,
int nCoeff,
int nSamples){
int nThreads = 32;
int nBlocks = nSamples/nThreads;
dim3 gridSz( nBlocks, 1, 1);
dim3 blockSz( nThreads, 1 , 1);
if( gpuBuf == NULL ) gpu_init();
70
CUDA_SAFE_CALL(
cudaMemcpy(gpuBuf, samples, sizeof(float)*(nSamples +
nCoeff), cudaMemcpyHostToDevice));
HPFilter<<< gridSz, blockSz >>>( coeff, gpuBuf, hpfBuf, nThreads, nBlocks );
CUDA_SAFE_CALL(
cudaMemcpy(out, hpfBuf, sizeof(float) * nSamples,
cudaMemcpyDeviceToHost)
);
}
/* Ajuste de escala de amostra */
__global__ void scaler_cuda(float *data_in,
float *data_out,
float scale){
int idx = (32 * blockIdx.x) + threadIdx.x;
data_out[idx] = data_in[idx] * scale;
}
void scaler(float *d,
int num_d,
float scl){
int nThreads = 32;
71
int nBlocks = num_d/nThreads;
if(num_d % nThreads) {
nBlocks++;
}
dim3 gridSz( nBlocks, 1, 1);
dim3 blockSz( nThreads, 1 , 1);
CUDA_SAFE_CALL(
cudaMemcpy(gpuBuf, d, sizeof(float)*num_d,
cudaMemcpyHostToDevice));
scaler_cuda<<< gridSz, blockSz >>>(gpuBuf, hpfBuf, scl);
CUDA_SAFE_CALL(
cudaMemcpy(d, hpfBuf, sizeof(float) * num_d,
cudaMemcpyDeviceToHost)
);
}
void scaler2(float *d,
float *d2,
int num_d,
float scl){
int nThreads = 32;
int nBlocks = (num_d * 2)/nThreads;
if((num_d *2) % nThreads) {
nBlocks++;
}
72
dim3 gridSz( nBlocks, 1, 1);
dim3 blockSz( nThreads, 1 , 1);
CUDA_SAFE_CALL(
cudaMemcpy(gpuBuf, d, sizeof(float)*num_d,
cudaMemcpyHostToDevice));
CUDA_SAFE_CALL(
cudaMemcpy(&gpuBuf[num_d], d2, sizeof(float)*num_d,
cudaMemcpyHostToDevice));
scaler_cuda<<< gridSz, blockSz >>>(gpuBuf, hpfBuf, scl);
CUDA_SAFE_CALL(
cudaMemcpy(d, hpfBuf, sizeof(float) * num_d,
cudaMemcpyDeviceToHost)
CUDA_SAFE_CALL(
cudaMemcpy(d2, &hpfBuf[num_d], sizeof(float) * num_d,
cudaMemcpyDeviceToHost)
}
);
);