Este Guia de Melhores Práticas é um manual para ajudar os desenvolvedores a obter o melhor desempenho das GPUs NVIDIA CUDA. Ele apresenta técnicas estabelecidas de paralelização e otimização e explica metáforas de codificação e idiomas que podem simplificar bastante a programação para arquiteturas de GPU compatíveis com CUDA. Enquanto o conteúdo pode ser usado como um manual de referência, você deve estar ciente de que alguns tópicos são revisados em diferentes contextos à medida que vários tópicos de programação e configuração são explorados. Como resultado, recomenda-se que os leitores da primeira vez passem pelo guia sequencialmente. Essa abordagem melhorará sua compreensão das práticas de programação efetivas e permitirá que você use melhor o guia para referência mais tarde. As discussões neste guia utilizam a linguagem de programação C, então você deve estar confortável ao ler o código C. Este guia refere-se a e depende de vários outros documentos que você deve ter à sua disposição para referência, todos os quais estão disponíveis sem nenhum custo do desenvolvedor do site CUDA. nvidiacuda-downloads. Os seguintes documentos são recursos especialmente importantes: Guia de Instalação do CUDA Guia de Programação do CUDA C Manual de Referência do CUDA Toolkit Em particular, a seção de otimização deste guia pressupõe que você já tenha baixado e instalado com sucesso o CUDA Toolkit (caso contrário, consulte o CUDA relevante Guia de instalação da sua plataforma) e que tenha familiaridade básica com a linguagem e o ambiente de programação do CUDA C (se não for, consulte o Guia de programação CUDA C). Este guia apresenta o ciclo de cálculo Avaliar, Paralelizar, Otimizar, Implantar (APOD) para aplicações com o objetivo de ajudar os desenvolvedores de aplicativos a identificar rapidamente as partes do seu código que se beneficiarão mais prontamente da aceleração GPU, perceberão rapidamente esse benefício e começarão a alavancar As acelerações resultantes em produção o mais cedo possível. O APOD é um processo cíclico: as acelerações iniciais podem ser alcançadas, testadas e implantadas com apenas um investimento inicial mínimo de tempo, momento em que o ciclo pode começar de novo, identificando novas oportunidades de otimização, ver velocidades adicionais e, em seguida, implementar as versões ainda mais rápidas do O aplicativo para a produção. Para um projeto existente, o primeiro passo é avaliar o aplicativo para localizar as partes do código que são responsáveis pela maior parte do tempo de execução. Armado com esse conhecimento, o desenvolvedor pode avaliar esses estrangulamentos para paralelização e começar a investigar a aceleração GPU. Ao entender os requisitos e restrições dos usuários finais e aplicando as leis Amdahls e Gustafsons, o desenvolvedor pode determinar o limite superior da melhoria do desempenho da aceleração das partes identificadas da aplicação. Tendo identificado os hotspots e tendo feito os exercícios básicos para estabelecer metas e expectativas, o desenvolvedor precisa paralelizar o código. Dependendo do código original, isso pode ser tão simples como chamar para uma biblioteca otimizada GPU existente, como o cuBLAS. QUFFT. Ou impulso. Ou pode ser tão simples como adicionar algumas diretivas de pré-processador como dicas para um compilador de paralelização. Por outro lado, alguns projetos de aplicativos exigirão alguma quantidade de refatoração para expor seu paralelismo inerente. Como mesmo futuras arquiteturas de CPU exigirão expor esse paralelismo para melhorar ou simplesmente manter o desempenho de aplicações seqüenciais, a família CUDA de linguagens de programação paralelas (CUDA CC, CUDA Fortran, etc.) tem como objetivo tornar a expressão desse paralelismo tão simples Possível, ao mesmo tempo que permite a operação em GPUs compatíveis com CUDA projetadas para o rendimento paralelo máximo. Depois de completar cada rodada de paralelização de aplicativos, o desenvolvedor pode se mover para otimizar a implementação para melhorar o desempenho. Uma vez que existem muitas otimizações possíveis que podem ser consideradas, ter uma boa compreensão das necessidades da aplicação pode ajudar a tornar o processo tão suave quanto possível. No entanto, como no APOD como um todo, a otimização de programas é um processo iterativo (identifique uma oportunidade de otimização, aplique e teste a otimização, verifique a aceleração alcançada e repita), o que significa que não é necessário que um programador gaste grandes quantidades De tempo memorizando a maior parte de todas as possíveis estratégias de otimização antes de ver boas acelerações. Em vez disso, as estratégias podem ser aplicadas de forma incremental à medida que são aprendidas. As otimizações podem ser aplicadas em vários níveis, desde a sobreposição de transferências de dados até a computação até o ajuste fino das seqüências de operação de ponto flutuante. As ferramentas de criação de perfil disponíveis são inestimáveis para orientar esse processo, pois podem ajudar a sugerir um próximo curso de ação para os esforços de otimização de desenvolvedores e fornecer referências nas partes relevantes da seção de otimização deste guia. Tendo completado a aceleração GPU de um ou mais componentes da aplicação, é possível comparar o resultado com a expectativa original. Lembre-se de que o passo de avaliação inicial permitiu ao desenvolvedor determinar um limite superior para a aceleração potencial alcançável através da aceleração de hotspots dados. Antes de abordar outros hotspots para melhorar a velocidade total, o desenvolvedor deve considerar levar a implementação parcialmente paralelizada e levá-la à produção. Isso é importante por vários motivos, por exemplo, permite ao usuário tirar proveito do investimento o mais cedo possível (a aceleração pode ser parcial, mas ainda é valiosa), e minimiza o risco para o desenvolvedor e o usuário, fornecendo uma evolução Em vez de um conjunto revolucionário de mudanças no aplicativo. Ao longo deste guia, são feitas recomendações específicas sobre a concepção e implementação do código CUDA C. Essas recomendações são categorizadas por prioridade, que é uma mistura do efeito da recomendação e seu alcance. As ações que apresentam melhorias substanciais para a maioria dos aplicativos CUDA têm a maior prioridade, enquanto pequenas otimizações que afetam apenas situações muito específicas recebem uma prioridade menor. Antes de implementar recomendações de menor prioridade, é uma boa prática certificar-se de que todas as recomendações de maior prioridade relevantes sejam aplicadas. Esta abordagem tenderá a fornecer os melhores resultados para o tempo investido e evitará a armadilha de otimização prematura. Os critérios de benefício e alcance para estabelecer a prioridade variam de acordo com a natureza do programa. Neste guia, eles representam um caso típico. Seu código pode refletir diferentes fatores de prioridade. Independentemente desta possibilidade, é uma boa prática verificar que nenhuma recomendação de maior prioridade tenha sido negligenciada antes de assumir itens de menor prioridade. Observação: as amostras de código ao longo do guia omitem a verificação de erros quanto a concisão. O código de produção deve, no entanto, verificar sistematicamente o código de erro retornado por cada chamada da API e verificar se há falhas nos lançamentos do kernel chamando cudaGetLastError (). Desde supercomputadores até telefones celulares, os processadores modernos dependem cada vez mais de paralelismo para oferecer desempenho. A unidade computacional central, que inclui controle, aritmética, registradores e normalmente algum cache, é replicada algumas vezes e conectada à memória através de uma rede. Como resultado, todos os processadores modernos requerem código paralelo para conseguir uma boa utilização do seu poder computacional. Enquanto os processadores estão evoluindo para expor mais paralelismo de grão fino ao programador, muitas aplicações existentes evoluíram como códigos em série ou como códigos paralelos de grão grosso (por exemplo, onde os dados são decompostos em regiões processadas em paralelo, com sub-regiões Compartilhado usando o MPI). Para se beneficiar de qualquer arquitetura de processador moderna, as GPUs incluídas, os primeiros passos são avaliar a aplicação para identificar os hotspots, determinar se eles podem ser paralelizados e entender as cargas de trabalho relevantes agora e no futuro. A programação do CUDA envolve a execução de código em duas plataformas diferentes simultaneamente: um sistema host com uma ou mais CPUs e um ou mais dispositivos GPU NVIDIA habilitados para CUDA. Enquanto as GPUs NVIDIA são freqüentemente associadas a gráficos, eles também são poderosos mecanismos aritméticos capazes de executar milhares de threads leves em paralelo. Essa capacidade os torna adequados para cálculos que podem alavancar a execução paralela. No entanto, o dispositivo é baseado em um design distintamente diferente do sistema host e é importante entender essas diferenças e como eles determinam o desempenho das aplicações CUDA para usar o CUDA efetivamente. As principais diferenças estão no modelo de threading e em memórias físicas separadas: Recursos de Threading As tubulações de execução em sistemas host podem suportar um número limitado de threads simultâneos. Os servidores que possuem quatro processadores hexadecimais hoje podem executar apenas 24 threads simultaneamente (ou 48 se as CPUs suportam o Hyper-Threading). Em comparação, a menor unidade executável de paralelismo em um dispositivo CUDA compreende 32 threads (denominada urdidura de threads) . As GPUs NVIDIA modernas podem suportar até 1536 threads ativas simultaneamente por multiprocessador (consulte Recursos e Especificações do Guia de Programação CUDA C) Em GPUs com 16 multiprocessadores, isso leva a mais de 24.000 threads ativos simultaneamente. Threads Threads em uma CPU geralmente são entidades de peso pesado. O sistema operacional deve trocar threads para dentro e fora dos canais de execução da CPU para fornecer capacidade de multithreading. Os interruptores de contexto (quando dois threads são trocados) são, portanto, lentos e caros. Em comparação, os tópicos em GPUs são extremamente leves. Em um sistema típico, milhares de threads são colocados em fila para o trabalho (em linhas de 32 threads cada). Se a GPU deve aguardar uma urdidura de threads, ela simplesmente começa a executar o trabalho em outra. Como os registros separados são alocados para todos os tópicos ativos, não é necessário trocar os registros ou outros estados ao mudar entre os segmentos da GPU. Os recursos permanecem alocados para cada segmento até que ele complete sua execução. Em suma, os núcleos da CPU são projetados para minimizar a latência para um ou dois segmentos ao mesmo tempo, enquanto as GPUs são projetadas para lidar com um grande número de threads simultâneos e leves para maximizar a taxa de transferência. RAM O sistema host e o dispositivo cada um têm suas próprias memórias físicas anexadas distintas. À medida que as memórias do host e do dispositivo são separadas pelo barramento PCI Express (PCIe), os itens na memória do host devem ocasionalmente ser comunicados através do barramento para a memória do dispositivo ou vice-versa, conforme descrito em O que funciona em um dispositivo habilitado para CUDA. Estes são os Diferenças primárias de hardware entre os hosts da CPU e os dispositivos GPU em relação à programação paralela. Outras diferenças são discutidas à medida que surgem em outros lugares neste documento. As aplicações compostas com estas diferenças em mente podem tratar o host e o dispositivo em conjunto como um sistema heterogêneo coesivo em que cada unidade de processamento é alavancada para fazer o tipo de trabalho que faz melhor: trabalho seqüencial no host e trabalho paralelo no dispositivo. Os seguintes problemas devem ser considerados ao determinar quais partes de um aplicativo executar no dispositivo: o dispositivo é ideal para cálculos que podem ser executados em vários elementos de dados simultaneamente em paralelo. Isso geralmente envolve aritmética em grandes conjuntos de dados (como matrizes), onde a mesma operação pode ser realizada em milhares, se não milhões, de elementos ao mesmo tempo. Este é um requisito para um bom desempenho no CUDA: o software deve usar um grande número (geralmente milhares ou dezenas de milhares) de threads simultâneos. O suporte para executar inúmeros fios em paralelo deriva do uso de CUDA de um modelo de encadeamento leve descrito acima. Para obter o melhor desempenho, deve haver alguma coerência no acesso à memória por threads adjacentes executados no dispositivo. Certos padrões de acesso à memória permitem que o hardware agrupe grupos de leituras ou gravações de vários itens de dados em uma única operação. Dados que não podem ser definidos de modo a permitir o coalescimento. Ou que não tem localidade suficiente para usar o L1 ou caches de textura de forma eficaz, tenderá a ver velocidades menores quando usado em cálculos no CUDA. Para usar CUDA, os valores de dados devem ser transferidos do host para o dispositivo ao longo do barramento PCI Express (PCIe). Essas transferências são caras em termos de desempenho e devem ser minimizadas. (Ver Transferência de Dados entre o Host eo Dispositivo.) Este custo tem várias ramificações: a complexidade das operações deve justificar o custo de mover dados de e para o dispositivo. O código que transfere dados para uso breve por um pequeno número de tópicos verá pouco ou nenhum benefício de desempenho. O cenário ideal é aquele em que muitos tópicos executam uma quantidade substancial de trabalho. Por exemplo, transferir duas matrizes para o dispositivo para executar uma adição de matriz e, em seguida, transferir os resultados para o host não perceberá muito benefício de desempenho. O problema aqui é o número de operações realizadas por elemento de dados transferido. Para o procedimento anterior, assumindo matrizes de tamanho NxN, há N 2 operações (adições) e 3N 2 elementos transferidos, de modo que a proporção de operações para elementos transferidos é 1: 3 ou O (1). Os benefícios de desempenho podem ser mais facilmente alcançados quando esta proporção é maior. Por exemplo, uma multiplicação de matriz das mesmas matrizes requer operações N3 (multiplicação), de modo que a relação das operações com os elementos transferidos é O (N), caso em que, quanto maior for a matriz, maior será o desempenho. Os tipos de operações são um fator adicional, pois as adições possuem diferentes perfis de complexidade do que, por exemplo, funções trigonométricas. É importante incluir a sobrecarga da transferência de dados para e do dispositivo para determinar se as operações devem ser realizadas no host ou no dispositivo. Os dados devem ser mantidos no dispositivo o maior tempo possível. Como as transferências devem ser minimizadas, os programas que executam vários kernels nos mesmos dados devem favorecer a saída dos dados no dispositivo entre as chamadas do kernel, em vez de transferir resultados intermediários para o host e enviá-los de volta ao dispositivo para cálculos subsequentes. Assim, no exemplo anterior, as duas matrizes a serem adicionadas já estavam no dispositivo como resultado de algum cálculo anterior, ou se os resultados da adição fossem usados em algum cálculo subseqüente, a adição da matriz deveria ser realizada localmente em o dispositivo. Essa abordagem deve ser usada mesmo que uma das etapas de uma seqüência de cálculos possa ser realizada mais rapidamente no host. Mesmo um kernel relativamente lento pode ser vantajoso se evita uma ou mais transferências PCIe. A Transferência de Dados entre o Host e o Dispositivo fornece mais detalhes, incluindo as medidas de largura de banda entre o host e o dispositivo em relação ao dispositivo propriamente dito. Em cache em L1 e L2 por padrão em dispositivos de capacidade de computação 2.x em cache somente em L2 por padrão em dispositivos de capacidades de computação mais altas, embora alguns permitam opt-in ao cache em L1 também através de sinalizadores de compilação. Em cache em L1 e L2 por padrão em dispositivos de capacidade de computação 2.x e 3.x dispositivos de capacidade de computação 5.x cache locals only in L2. No caso do acesso à textura, se uma referência de textura for vinculada a uma matriz linear na memória global, o código do dispositivo pode escrever na matriz subjacente. As referências de textura que são vinculadas às matrizes CUDA podem ser escritas através de operações de escrita de superfície, vinculando uma superfície ao mesmo armazenamento de matriz CUDA subjacente). A leitura de uma textura enquanto se escreve para a matriz de memória global subjacente no mesmo lançamento do kernel deve ser evitada porque os caches de textura são somente leitura e não são invalidados quando a memória global associada é modificada. Talvez a consideração de desempenho mais importante na programação para arquiteturas de GPU compatíveis com CUDA seja o coalescimento de acessos de memória global. As cargas de memória global e as lojas por threads de uma urdidura são combinadas pelo dispositivo em apenas uma transação quando determinados requisitos de acesso são atendidos. Nota: Alta prioridade: Assegure-se de que os acessos globais de memória se agrupam sempre que possível. Os requisitos de acesso para coalescer dependem da capacidade de cálculo do dispositivo e estão documentados no Guia de Programação CUDA C. Para dispositivos de capacidade de computação 2.x, os requisitos podem ser resumidos com bastante facilidade: os acessos simultâneos dos segmentos de uma urdidura se juntarão em várias transações iguais ao número de linhas de cache necessárias para atender todos os segmentos da urdidura . Por padrão, todos os acessos são armazenados em cache por L1, que como linhas de 128 bytes. Para padrões de acesso dispersos, para reduzir o excesso de estoque, às vezes pode ser útil armazenar em cache apenas em L2, que armazena segmentos de 32 bytes mais curtos (consulte o Guia de Programação CUDA C). Para dispositivos de capacidade de computação 3.x, os acessos a memória global são armazenados em cache somente em L2 L1 é reservado para acessos de memória local. Alguns dispositivos de capacidade de computação 3.5, 3.7 ou 5.2 permitem o armazenamento em cache opt-in de globals em L1 também. O acesso à memória de forma conjunta é ainda mais importante quando o ECC está ligado. Os acessos dispersos aumentam a sobrecarga de transferência de memória ECC, especialmente quando se escrevem dados para a memória global. Os conceitos de coalescência são ilustrados nos seguintes exemplos simples. Estes exemplos assumem a capacidade de cálculo 2.x. Esses exemplos assumem que os acessos são armazenados em cache através de L1, que é o comportamento padrão desses dispositivos e que os acessos são para palavras de 4 bytes, a menos que seja observado de outra forma. O primeiro e mais simples caso de coalescência pode ser alcançado por qualquer dispositivo habilitado para CUDA: o segmento k-th acessa a palavra k-ésima em uma linha de cache. Nem todos os tópicos precisam participar. Por exemplo, se os segmentos de um acesso de urdidura adjacente a palavras de 4 bytes (por exemplo, valores de flutuação adjacentes), uma única linha de cache L1 de 128B e, portanto, uma única transação coalescida servirá esse acesso à memória. Esse padrão é mostrado na Figura 3. Figura 3. Acesso coalescido - todos os segmentos acessam uma linha de cache Este padrão de acesso resulta em uma única transação L1 de 128 bytes, indicada pelo retângulo vermelho. Se algumas palavras da linha não haviam sido solicitadas por nenhum tópico (como se vários segmentos tivessem acessado a mesma palavra ou se alguns tópicos não participaram do acesso), todos os dados na linha de cache são obtidos de qualquer maneira. Além disso, se os acessos pelos tópicos da urdidura fossem permutados dentro desse segmento, ainda uma transação L1 de 128 bytes teria sido executada por um dispositivo com capacidade de cálculo 2.x. Se os segmentos seqüenciais em uma memória de acesso de urdidura que é seqüencial, mas não alinhados com as linhas de cache, serão solicitadas duas caché L1 de 128 bytes, como mostrado na Figura 4. Figura 4. Endereços sequenciais não alinhados que se encaixam em dois 128-bytes L1- Linhas de cache Para transações que não sejam de cache (ou seja, aqueles que ignoram L1 e usam apenas o cache L2), um efeito semelhante é visto, exceto no nível dos segmentos de L2 de 32 bytes. Na Figura 5., vemos um exemplo disso: o mesmo padrão de acesso da Figura 4 é usado, mas agora o cache L1 está desativado, então agora são necessários 5 segmentos L2 de 32 bytes para satisfazer a solicitação. Figura 5. Endereços seqüenciais desalinhados que se enquadram em cinco segmentos de cache L2 de 32 bytes Memória alocada através da API do Runtime do CUDA, como via cudaMalloc (). É garantido para ser alinhado a pelo menos 256 bytes. Portanto, a escolha de tamanhos de bloco de thread sensíveis, como múltiplos do tamanho da urdidura (ou seja, 32 em GPUs atuais), facilita os acessos de memória por urdiduras alinhadas às linhas de cache. (Considere o que aconteceria com os endereços de memória acessados pelo segundo, terceiro e subseqüentes blocos de thread se o tamanho do bloco de thread não fosse um múltiplo do tamanho da urdidura, por exemplo.) É fácil e informativo explorar as ramificações de acessos desalinhados usando Um kernel de cópia simples, como aquele em um kernel de cópia A que ilustra acessos desalinhados. Um kernel de cópia que ilustra acessos desalinhados em um kernel de cópia que ilustra acessos desalinhados. Os dados são copiados da matriz de entrada idata para a matriz de saída, ambas existentes na memória global. O kernel é executado dentro de um loop no código do host que varia o deslocamento do parâmetro de 0 a 32. (A Figura 4 e a Figura 4 correspondem a desalinhamentos nos casos de caching e acessos de memória sem cache, respectivamente.) A largura de banda efetiva para a cópia Com várias compensações em uma NVIDIA Tesla M2090 (capacidade de cálculo 2.0, com ECC ativado, como é por padrão) é mostrado na Figura 6. Figura 6. Desempenho do kernel offsetCopy Para a NVIDIA Tesla M2090, a memória global acessa sem deslocamento ou Com offsets que são múltiplos de 32 palavras resultam em uma única transação de linha de cache L1 ou 4 cargas de segmento de cache L2 (para cargas de cache não L1). A largura de banda alcançada é de aproximadamente 130GBs. Caso contrário, duas linhas de cache L1 (modo de cache) ou quatro a cinco segmentos de cache L2 (modo sem cache) são carregadas por urdidura, resultando em aproximadamente 45º da taxa de transferência de memória alcançada sem deslocamentos. Um ponto interessante é que podemos esperar que o caso de cache seja mais pior do que o caso de não armazenamento em cache para essa amostra, dado que cada urdidura no caso de armazenamento em cache obtém o dobro de bytes necessários, enquanto que no caso de não cache, apenas 54 tantos bytes quanto necessário são obtidos por urdidura. Neste exemplo particular, esse efeito não é aparente, no entanto, porque as urdiduras adjacentes reutilizam as linhas de cache que seus vizinhos buscaram. Então, enquanto o impacto ainda é evidente no caso de colocação em cache de cargas, não é tão bom quanto esperávamos. Teria sido muito mais, se as urdiduras adjacentes não exibissem um alto grau de reutilização das linhas de cache excedentes. Como visto acima, no caso de acessos seqüenciais desalinhados, os caches dos dispositivos 2.x da computação ajudam muito a alcançar um desempenho razoável. Pode ser diferente com acessos não rígidos, no entanto, e esse é um padrão que ocorre com freqüência quando se trata de dados ou matrizes multidimensionais. Por esse motivo, garantir que, tanto quanto possível, os dados em cada linha de cache obtida seja realmente usada é uma parte importante da otimização de desempenho de acessos de memória nesses dispositivos. Para ilustrar o efeito do acesso strided na largura de banda efetiva, veja o kernel strideCopy () no kernel de A para ilustrar a cópia de dados de stride não-unidade. Que copia dados com passos de stride entre threads de idata para odata. Um kernel para ilustrar a cópia de dados de stride não-unidade Figura 7 ilustra essa situação neste caso, faz threads dentro de uma palavra de acesso à urdidura na memória com um passo de 2. Essa ação leva a uma carga de duas linhas de cache L1 (ou cache de oito L2 Segmentos no modo não cache) por urdidura no Tesla M2090 (capacidade de cálculo 2.0). Figura 7. Roscas adjacentes acessando a memória com um passo de 2 Um passo de 2 resultados em 50 de eficiência de armazenamento de carga, pois metade dos elementos na transação não são usados e representam a largura de banda desperdiçada. À medida que o stride aumenta, a largura de banda efetiva diminui até o ponto em que 32 linhas de cache são carregadas para os 32 threads em uma urdidura, conforme indicado na Figura 8. Figura 8. Desempenho do kernel strideCopy Conforme ilustrado na Figura 8. não-unidade - Os acessos de memória globais devem ser evitados sempre que possível. Um método para isso usa memória compartilhada, que é discutida na próxima seção. Como é no chip, a memória compartilhada possui largura de banda muito maior e menor latência que a memória local e global - desde que não haja conflitos bancários entre os segmentos, conforme detalhado na seção a seguir. Para obter alta largura de banda de memória para acessos simultâneos, a memória compartilhada é dividida em módulos de memória de tamanho igual (bancos) que podem ser acessados simultaneamente. Portanto, qualquer carga de memória ou armazenamento de n endereços que abrange n bancos de memória distintos pode ser atendida simultaneamente, produzindo uma largura de banda efetiva que é n vezes maior que a largura de banda de um único banco. No entanto, se vários endereços de um mapa de solicitação de memória para o mesmo banco de memória, os acessos são serializados. O hardware divide uma solicitação de memória que tem conflitos bancários em tantos pedidos separados sem conflitos quanto necessário, diminuindo a largura de banda efetiva por um fator igual ao número de solicitações de memória separadas. A única exceção aqui é quando vários tópicos em um endereço de urdidura são a mesma localização de memória compartilhada, resultando em uma transmissão. Os dispositivos da capacidade de computação 2.x e superior têm a capacidade adicional de acessos de memória compartilhada multicast (ou seja, enviar cópias do mesmo valor para vários segmentos da urdidura). Para minimizar os conflitos bancários, é importante entender como os endereços de memória são mapeados para bancos de memória e como agendar de forma otimizada pedidos de memória. Compute Capability 2.x Em dispositivos de capacidade de computação 2.x, cada banco possui uma largura de banda de 32 bits a cada dois ciclos de clock e sucessivas palavras de 32 bits são atribuídas a bancos sucessivos. O tamanho da urdidura é de 32 threads e o número de bancos também é 32, portanto, conflitos bancários podem ocorrer entre qualquer segmento na urdidura. Consulte Calcular capacidade 2.x no Guia de programação CUDA C para obter mais detalhes. Compute Capability 3.x Em dispositivos de capacidade de computação 3.x, cada banco possui uma largura de banda de 64 bits a cada ciclo de clock (). Existem dois modos bancários diferentes: as palavras sucessivas de 32 bits (no modo de 32 bits) ou sucessivas palavras de 64 bits (modo de 64 bits) são atribuídas a bancos sucessivos. O tamanho da urdidura é de 32 threads e o número de bancos também é 32, portanto, conflitos bancários podem ocorrer entre qualquer segmento na urdidura. Consulte Compute Capability 3.x no CUDA C Programming Guide para obter mais detalhes. Nota: () No entanto, os dispositivos de capacidade de computação 3.x normalmente têm freqüências de clock mais baixas do que os dispositivos de capacidade de computação 2.x para melhorar a eficiência de energia. A memória compartilhada permite a cooperação entre threads em um bloco. Quando vários threads em um bloco usam os mesmos dados da memória global, a memória compartilhada pode ser usada para acessar os dados da memória global apenas uma vez. A memória compartilhada também pode ser usada para evitar acessos de memória não vendidos carregando e armazenando dados em um padrão coalescido da memória global e depois reordenando-o na memória compartilhada. Além dos conflitos de banco de memória, não há penalidade para acessos não sequenciais ou não alinhados por uma urdidura na memória compartilhada. O uso da memória compartilhada é ilustrado através do exemplo simples de uma multiplicação de matriz C AB para o caso com A de dimensão Mxw, B de dimensão wxN e C de dimensão MxN. Para manter os kernels simples, M e N são múltiplos de 32, e w é 32 para dispositivos de capacidade de computação 2.0 ou superior. Uma decomposição natural do problema é usar um tamanho de bloco e telha de fios wxw. Portanto, em termos de telhas wxw, A é uma matriz de coluna, B é uma matriz de linha e C é o produto externo, veja a Figura 9. É iniciada uma grade de Nw por blocos Mw, onde cada bloco de segmento calcula os elementos de um diferente Azulejo em C de uma única tela de A e uma única tela de B. Figura 9. Matriz de bloco-coluna multiplicada por matriz de bloco-linha. Matriz de bloco-coluna (A) multiplicada pela matriz de bloco-linha (B) com matriz de produto resultante (C). Para fazer isso, o núcleo simples do Múltiplo (multiplicação da matriz não optimizada) calcula os elementos de saída de um mosaico da matriz C. Multiplicação da matriz não optimizada Na multiplicação da matriz não optimizada. uma. B. E c são ponteiros para a memória global para as matrizes A, B e C, respectivamente, blockDim. x. BlockDim. y. E TILEDIM são todos iguais a w. Cada segmento no bloco wxw-thread calcula um elemento em uma tela de C. row e col são a linha e a coluna do elemento em C sendo calculada por um segmento particular. O loop for over multiplica uma linha de A por uma coluna de B, que é então escrita em C. A largura de banda efetiva deste kernel é de apenas 6,6 GBs em um NVIDIA Tesla K20X (com ECC desativado). Para analisar o desempenho, é necessário considerar como os warps acessam a memória global no loop for. Cada urdidura de threads calcula uma linha de uma telha de C, que depende de uma única linha de A e de uma telha inteira de B como ilustrado na Figura 10. Figura 10. Computação de uma linha de uma telha. Computação de uma linha de uma telha em C usando uma linha de A e uma telha inteira de B. Para cada iteração i do loop for, os threads em uma urdidura lêem uma linha do azulejo B, que é um acesso seqüencial e coalescido para Todos os recursos de computação. No entanto, para cada iteração i. Todos os tópicos em uma urdidura lêem o mesmo valor da memória global para a matriz A, pois o índice rowTILEDIMi é constante dentro de uma urdidura. Mesmo que esse acesso requer apenas 1 transação em dispositivos de capacidade de computação 2.0 ou superior, há uma largura de banda desperdiçada na transação, porque apenas uma palavra de 4 bytes fora de 32 palavras na linha de cache é usada. Podemos reutilizar esta linha de cache em iterações subseqüentes do loop, e, eventualmente, utilizamos todas as 32 palavras, no entanto, quando muitas veias executam no mesmo multiprocessador simultaneamente, como é geralmente o caso, a linha de cache pode ser facilmente despejada do cache entre Iterações i e i1. O desempenho em um dispositivo de qualquer capacidade de computação pode ser melhorado lendo um mosaico de A em memória compartilhada como mostrado em Usando memória compartilhada para melhorar a eficiência de carga de memória global na multiplicação de matriz. Usando memória compartilhada para melhorar a eficiência de carga de memória global na multiplicação de matriz. Usando memória compartilhada para melhorar a eficiência de carga de memória global na multiplicação de matriz. Cada elemento em uma telha de A é lido da memória global apenas uma vez, de forma totalmente compartilhada (sem largura de banda desperdiçada), para a memória compartilhada. Dentro de cada iteração do loop for, um valor na memória compartilhada é transmitido para todos os tópicos em uma urdidura. Não é necessária nenhuma chamada de barreira de sincronização syncthreads (), depois de ler a telha de A na memória compartilhada porque somente os segmentos dentro da urdidura que escrevem os dados na memória compartilhada lêem os dados (Nota: em vez de syncthreads (). A matriz compartilhada pode precisar Seja marcado como volátil para a correção em dispositivos de capacidade de computação 2.0 ou superior, veja o Guia de Compatibilidade NVIDIA Fermi). Este kernel possui uma largura de banda efetiva de 7.8GBs em um NVIDIA Tesla K20X. Isso ilustra o uso da memória compartilhada como um cache gerenciado pelo usuário quando a política de despejo de cache L1 do hardware não coincide bem com as necessidades do aplicativo ou quando o cache L1 não é usado para ler da memória global. Uma melhoria adicional pode ser feita sobre como usar a memória compartilhada para melhorar a eficiência de carga de memória global na multiplicação de matriz trata da matriz B. Ao calcular cada uma das linhas de uma telha da matriz C, lê toda a telha de B. A leitura repetida do azulejo B pode ser eliminada lendo-o na memória compartilhada uma vez (Melhoria lendo dados adicionais na memória compartilhada). Melhoria ao ler dados adicionais na memória compartilhada. Observe que, em Melhoria, lendo dados adicionais na memória compartilhada. Uma ligação syncthreads () é necessária depois de ler o mosaico B porque uma urdidura lê dados de memória compartilhada que foram escritos em memória compartilhada por warps diferentes. The effective bandwidth of this routine is 14.9 GBs on an NVIDIA Tesla K20X. Note that the performance improvement is not due to improved coalescing in either case, but to avoiding redundant transfers from global memory. The results of the various optimizations are summarized in Table 2. Table 2. Performance Improvements Optimizing C AB Matrix Multiply NVIDIA Tesla K20X These results should be compared with those in Table 2. As can be seen from these tables, judicious use of shared memory can dramatically improve performance. The examples in this section have illustrated three reasons to use shared memory: To enable coalesced accesses to global memory, especially to avoid large strides (for general matrices, strides are much larger than 32) To eliminate (or reduce) redundant loads from global memory To avoid wasted bandwidth Local memory is so named because its scope is local to the thread, not because of its physical location. In fact, local memory is off-chip. Hence, access to local memory is as expensive as access to global memory. In other words, the term local in the name does not imply faster access. Local memory is used only to hold automatic variables. This is done by the nvcc compiler when it determines that there is insufficient register space to hold the variable. Automatic variables that are likely to be placed in local memory are large structures or arrays that would consume too much register space and arrays that the compiler determines may be indexed dynamically. Inspection of the PTX assembly code (obtained by compiling with - ptx or - keep command-line options to nvcc ) reveals whether a variable has been placed in local memory during the first compilation phases. If it has, it will be declared using the. local mnemonic and accessed using the ld. local and st. local mnemonics. If it has not, subsequent compilation phases might still decide otherwise, if they find the variable consumes too much register space for the targeted architecture. There is no way to check this for a specific variable, but the compiler reports total local memory usage per kernel (lmem) when run with the --ptxas-options-v option. The read-only texture memory space is cached. Therefore, a texture fetch costs one device memory read only on a cache miss otherwise, it just costs one read from the texture cache. The texture cache is optimized for 2D spatial locality, so threads of the same warp that read texture addresses that are close together will achieve best performance. Texture memory is also designed for streaming fetches with a constant latency that is, a cache hit reduces DRAM bandwidth demand, but not fetch latency. In certain addressing situations, reading device memory through texture fetching can be an advantageous alternative to reading device memory from global or constant memory. If textures are fetched using tex1D(). tex2D(). or tex3D() rather than tex1Dfetch(). the hardware provides other capabilities that might be useful for some applications such as image processing, as shown in Table 4. Table 4. Useful Features for tex1D(), tex2D(), and tex3D() Fetches Automatic handling of boundary cases 1 Can be used only with normalized texture coordinates 1 The automatic handling of boundary cases in the bottom row of Table 4 refers to how a texture coordinate is resolved when it falls outside the valid addressing range. There are two options: clamp and wrap. If x is the coordinate and N is the number of texels for a one-dimensional texture, then with clamp, x is replaced by 0 if x lt 0 and by 1-1 N if 1 lt x . With wrap, x is replaced by frac(x) where frac(x) x - floor(x) . Floor returns the largest integer less than or equal to x . So, in clamp mode where N 1, an x of 1.3 is clamped to 1.0 whereas in wrap mode, it is converted to 0.3 Within a kernel call, the texture cache is not kept coherent with respect to global memory writes, so texture fetches from addresses that have been written via global stores in the same kernel call return undefined data. That is, a thread can safely read a memory location via texture if the location has been updated by a previous kernel call or memory copy, but not if it has been previously updated by the same thread or another thread within the same kernel call. There is a total of 64 KB constant memory on a device. The constant memory space is cached. As a result, a read from constant memory costs one memory read from device memory only on a cache miss otherwise, it just costs one read from the constant cache. Accesses to different addresses by threads within a warp are serialized, thus the cost scales linearly with the number of unique addresses read by all threads within a warp. As such, the constant cache is best when threads in the same warp accesses only a few distinct locations. If all threads of a warp access the same location, then constant memory can be as fast as a register access. Generally, accessing a register consumes zero extra clock cycles per instruction, but delays may occur due to register read-after-write dependencies and register memory bank conflicts. The latency of read-after-write dependencies is approximately 24 cycles, but this latency is completely hidden on multiprocessors that have sufficient warps of threads concurrent per multiprocessor. For devices of compute capability 2.0, which have 32 CUDA cores per multiprocessor, as many as 768 threads (24 warps) might be required to completely hide latency, and so on for devices of higher compute capabilities. The compiler and hardware thread scheduler will schedule instructions as optimally as possible to avoid register memory bank conflicts. They achieve the best results when the number of threads per block is a multiple of 64. Other than following this rule, an application has no direct control over these bank conflicts. In particular, there is no register-related reason to pack data into float4 or int4 types. Register pressure occurs when there are not enough registers available for a given task. Even though each multiprocessor contains thousands of 32-bit registers (see Features and Technical Specifications of the CUDA C Programming Guide ), these are partitioned among concurrent threads. To prevent the compiler from allocating too many registers, use the - maxrregcountN compiler command-line option (see nvcc ) or the launch bounds kernel definition qualifier (see Execution Configuration of the CUDA C Programming Guide ) to control the maximum number of registers to allocated per thread. Device memory allocation and de-allocation via cudaMalloc() and cudaFree() are expensive operations, so device memory should be reused andor sub-allocated by the application wherever possible to minimize the impact of allocations on overall performance. One of the keys to good performance is to keep the multiprocessors on the device as busy as possible. A device in which work is poorly balanced across the multiprocessors will deliver suboptimal performance. Hence, its important to design your application to use threads and blocks in a way that maximizes hardware utilization and to limit practices that impede the free distribution of work. A key concept in this effort is occupancy, which is explained in the following sections. Hardware utilization can also be improved in some cases by designing your application so that multiple, independent kernels can execute at the same time. Multiple kernels executing at the same time is known as concurrent kernel execution. Concurrent kernel execution is described below. Another important concept is the management of system resources allocated for a particular task. How to manage this resource utilization is discussed in the final sections of this chapter. Thread instructions are executed sequentially in CUDA, and, as a result, executing other warps when one warp is paused or stalled is the only way to hide latencies and keep the hardware busy. Some metric related to the number of active warps on a multiprocessor is therefore important in determining how effectively the hardware is kept busy. This metric is occupancy. Occupancy is the ratio of the number of active warps per multiprocessor to the maximum number of possible active warps. (To determine the latter number, see the deviceQuery CUDA Sample or refer to Compute Capabilities in the CUDA C Programming Guide .) Another way to view occupancy is the percentage of the hardwares ability to process warps that is actively in use. Higher occupancy does not always equate to higher performance-there is a point above which additional occupancy does not improve performance. However, low occupancy always interferes with the ability to hide memory latency, resulting in performance degradation. One of several factors that determine occupancy is register availability. Register storage enables threads to keep local variables nearby for low-latency access. However, the set of registers (known as the register file ) is a limited commodity that all threads resident on a multiprocessor must share. Registers are allocated to an entire block all at once. So, if each thread block uses many registers, the number of thread blocks that can be resident on a multiprocessor is reduced, thereby lowering the occupancy of the multiprocessor. The maximum number of registers per thread can be set manually at compilation time per-file using the - maxrregcount option or per-kernel using the launchbounds qualifier (see Register Pressure ). For purposes of calculating occupancy, the number of registers used by each thread is one of the key factors. For example, devices with compute capability 1.1 have 8,192 32-bit registers per multiprocessor and can have a maximum of 768 simultaneous threads resident (24 warps x 32 threads per warp). This means that in one of these devices, for a multiprocessor to have 100 occupancy, each thread can use at most 10 registers. However, this approach of determining how register count affects occupancy does not take into account the register allocation granularity. For example, on a device of compute capability 1.1, a kernel with 128-thread blocks using 12 registers per thread results in an occupancy of 83 with 5 active 128-thread blocks per multi-processor, whereas a kernel with 256-thread blocks using the same 12 registers per thread results in an occupancy of 66 because only two 256-thread blocks can reside on a multiprocessor. Furthermore, register allocations are rounded up to the nearest 256 registers per block on devices with compute capability 1.1. The number of registers available, the maximum number of simultaneous threads resident on each multiprocessor, and the register allocation granularity vary over different compute capabilities. Because of these nuances in register allocation and the fact that a multiprocessors shared memory is also partitioned between resident thread blocks, the exact relationship between register usage and occupancy can be difficult to determine. The --ptxas optionsv option of nvcc details the number of registers used per thread for each kernel. See Hardware Multithreading of the CUDA C Programming Guide for the register allocation formulas for devices of various compute capabilities and Features and Technical Specifications of the CUDA C Programming Guide for the total number of registers available on those devices. Alternatively, NVIDIA provides an occupancy calculator in the form of an Excel spreadsheet that enables developers to hone in on the optimal balance and to test different possible scenarios more easily. This spreadsheet, shown in Figure 11. is called CUDAOccupancyCalculator. xls and is located in the tools subdirectory of the CUDA Toolkit installation. Figure 11. Using the CUDA Occupancy Calculator to project GPU multiprocessor occupancy In addition to the calculator spreadsheet, occupancy can be determined using the NVIDIA Visual Profilers Achieved Occupancy metric. The Visual Profiler also calculates occupancy as part of the Multiprocessor stage of application analysis. As described in Asynchronous and Overlapping Transfers with Computation. CUDA streams can be used to overlap kernel execution with data transfers. On devices that are capable of concurrent kernel execution, streams can also be used to execute multiple kernels simultaneously to more fully take advantage of the devices multiprocessors. Whether a device has this capability is indicated by the concurrentKernels field of the cudaDeviceProp structure (or listed in the output of the deviceQuery CUDA Sample). Non-default streams (streams other than stream 0) are required for concurrent execution because kernel calls that use the default stream begin only after all preceding calls on the device (in any stream) have completed, and no operation on the device (in any stream) commences until they are finished. The following example illustrates the basic technique. Because kernel1 and kernel2 are executed in different, non-default streams, a capable device can execute the kernels at the same time. CUDA work occurs within a process space for a particular GPU known as a context. The context encapsulates kernel launches and memory allocations for that GPU as well as supporting constructs such as the page tables. The context is explicit in the CUDA Driver API but is entirely implicit in the CUDA Runtime API, which creates and manages contexts automatically. With the CUDA Driver API, a CUDA application process can potentially create more than one context for a given GPU. If multiple CUDA application processes access the same GPU concurrently, this almost always implies multiple contexts, since a context is tied to a particular host process unless CUDA Multi-Process Service is in use. While multiple contexts (and their associated resources such as global memory allocations) can be allocated concurrently on a given GPU, only one of these contexts can execute work at any given moment on that GPU contexts sharing the same GPU are time-sliced. Creating additional contexts incurs memory overhead for per-context data and time overhead for context switching. Furthermore, the need for context switching can reduce utilization when work from several contexts could otherwise execute concurrently (see also Concurrent Kernel Execution ). Therefore, it is best to avoid multiple contexts per GPU within the same CUDA application. To assist with this, the CUDA Driver API provides methods to access and manage a special context on each GPU called the primary context. These are the same contexts used implicitly by the CUDA Runtime when there is not already a current context for a thread. Note: NVIDIA-SMI can be used to configure a GPU for exclusive compute mode. which limits the number of threads andor processes that can have simultaneous contexts on a particular GPU to one. Note: Medium Priority: To hide latency arising from register dependencies, maintain sufficient numbers of active threads per multiprocessor (i. e. sufficient occupancy). Register dependencies arise when an instruction uses a result stored in a register written by an instruction before it. The latency on current CUDA-enabled GPUs is approximately 24 cycles, so threads must wait 24 cycles before using an arithmetic result. However, this latency can be completely hidden by the execution of threads in other warps. See Registers for details. Note: Medium Priority: The number of threads per block should be a multiple of 32 threads, because this provides optimal computing efficiency and facilitates coalescing. The dimension and size of blocks per grid and the dimension and size of threads per block are both important factors. The multidimensional aspect of these parameters allows easier mapping of multidimensional problems to CUDA and does not play a role in performance. As a result, this section discusses size but not dimension. Latency hiding and occupancy depend on the number of active warps per multiprocessor, which is implicitly determined by the execution parameters along with resource (register and shared memory) constraints. Choosing execution parameters is a matter of striking a balance between latency hiding (occupancy) and resource utilization. Choosing the execution configuration parameters should be done in tandem however, there are certain heuristics that apply to each parameter individually. When choosing the first execution configuration parameter-the number of blocks per grid, or grid size - the primary concern is keeping the entire GPU busy. The number of blocks in a grid should be larger than the number of multiprocessors so that all multiprocessors have at least one block to execute. Furthermore, there should be multiple active blocks per multiprocessor so that blocks that arent waiting for a syncthreads() can keep the hardware busy. This recommendation is subject to resource availability therefore, it should be determined in the context of the second execution parameter - the number of threads per block, or block size - as well as shared memory usage. To scale to future devices, the number of blocks per kernel launch should be in the thousands. When choosing the block size, it is important to remember that multiple concurrent blocks can reside on a multiprocessor, so occupancy is not determined by block size alone. In particular, a larger block size does not imply a higher occupancy. For example, on a device of compute capability 1.1 or lower, a kernel with a maximum block size of 512 threads results in an occupancy of 66 percent because the maximum number of threads per multiprocessor on such a device is 768. Hence, only a single block can be active per multiprocessor. However, a kernel with 256 threads per block on such a device can result in 100 percent occupancy with three resident active blocks. As mentioned in Occupancy. higher occupancy does not always equate to better performance. For example, improving occupancy from 66 percent to 100 percent generally does not translate to a similar increase in performance. A lower occupancy kernel will have more registers available per thread than a higher occupancy kernel, which may result in less register spilling to local memory. Typically, once an occupancy of 50 percent has been reached, additional increases in occupancy do not translate into improved performance. It is in some cases possible to fully cover latency with even fewer warps, notably via instruction-level parallelism (ILP) for discussion, see nvidiacontentGTC-2010pdfs2238GTC2010.pdf. There are many such factors involved in selecting block size, and inevitably some experimentation is required. However, a few rules of thumb should be followed: Threads per block should be a multiple of warp size to avoid wasting computation on under-populated warps and to facilitate coalescing. A minimum of 64 threads per block should be used, and only if there are multiple concurrent blocks per multiprocessor. Between 128 and 256 threads per block is a better choice and a good initial range for experimentation with different block sizes. Use several (3 to 4) smaller thread blocks rather than one large thread block per multiprocessor if latency affects performance. This is particularly beneficial to kernels that frequently call syncthreads(). Note that when a thread block allocates more registers than are available on a multiprocessor, the kernel launch fails, as it will when too much shared memory or too many threads are requested. Shared memory can be helpful in several situations, such as helping to coalesce or eliminate redundant access to global memory. However, it also can act as a constraint on occupancy. In many cases, the amount of shared memory required by a kernel is related to the block size that was chosen, but the mapping of threads to shared memory elements does not need to be one-to-one. For example, it may be desirable to use a 32x32 element shared memory array in a kernel, but because the maximum number of threads per block is 512, it is not possible to launch a kernel with 32x32 threads per block. In such cases, kernels with 32x16 or 32x8 threads can be launched with each thread processing two or four elements, respectively, of the shared memory array. The approach of using a single thread to process multiple elements of a shared memory array can be beneficial even if limits such as threads per block are not an issue. This is because some operations common to each element can be performed by the thread once, amortizing the cost over the number of shared memory elements processed by the thread. A useful technique to determine the sensitivity of performance to occupancy is through experimentation with the amount of dynamically allocated shared memory, as specified in the third parameter of the execution configuration. By simply increasing this parameter (without modifying the kernel), it is possible to effectively reduce the occupancy of the kernel and measure its effect on performance. As mentioned in the previous section, once an occupancy of more than 50 percent has been reached, it generally does not pay to optimize parameters to obtain higher occupancy ratios. The previous technique can be used to determine whether such a plateau has been reached. Awareness of how instructions are executed often permits low-level optimizations that can be useful, especially in code that is run frequently (the so-called hot spot in a program). Best practices suggest that this optimization be performed after all higher-level optimizations have been completed. Single-precision floats provide the best performance, and their use is highly encouraged. The throughput of individual arithmetic operations is detailed in the CUDA C Programming Guide. Note: Low Priority: Use shift operations to avoid expensive division and modulo calculations. Integer division and modulo operations are particularly costly and should be avoided or replaced with bitwise operations whenever possible: If n is a power of 2, ( i n ) is equivalent to ( i log2 n ) and ( i n ) is equivalent to ( i amp n - 1 ). The compiler will perform these conversions if n is literal. (For further information, refer to Performance Guidelines in the CUDA C Programming Guide ). The reciprocal square root should always be invoked explicitly as rsqrtf() for single precision and rsqrt() for double precision. The compiler optimizes 1.0fsqrtf(x) into rsqrtf() only when this does not violate IEEE-754 semantics. Note: Low Priority: Avoid automatic conversion of doubles to floats. The compiler must on occasion insert conversion instructions, introducing additional execution cycles. This is the case for: Functions operating on char or short whose operands generally need to be converted to an int Double-precision floating-point constants (defined without any type suffix) used as input to single-precision floating-point computations The latter case can be avoided by using single-precision floating-point constants, defined with an f suffix such as 3.141592653589793f. 1.0f. 0.5f. This suffix has accuracy implications in addition to its ramifications on performance. The effects on accuracy are discussed in Promotions to Doubles and Truncations to Floats. Note that this distinction is particularly important to performance on devices of compute capability 2.x. For single-precision code, use of the float type and the single-precision math functions are highly recommended. When compiling for devices without native double-precision support such as devices of compute capability 1.2 and earlier, each double-precision floating-point variable is converted to single-precision floating-point format (but retains its size of 64 bits) and double-precision arithmetic is demoted to single-precision arithmetic. It should also be noted that the CUDA math librarys complementary error function, erfcf(). is particularly fast with full single-precision accuracy. For some fractional exponents, exponentiation can be accelerated significantly compared to the use of pow() by using square roots, cube roots, and their inverses. For those exponentiations where the exponent is not exactly representable as a floating-point number, such as 13, this can also provide much more accurate results, as use of pow() magnifies the initial representational error. The formulas in the table below are valid for x gt 0, x -0. that is, signbit(x) 0. Table 5. Formulae for exponentiation by small fractions Note: Medium Priority: Use the fast math library whenever speed trumps precision. Two types of runtime math operations are supported. They can be distinguished by their names: some have names with prepended underscores, whereas others do not (e. g. functionName() versus functionName() ). Functions following the functionName() naming convention map directly to the hardware level. They are faster but provide somewhat lower accuracy (e. g. sinf(x) and expf(x) ). Functions following functionName() naming convention are slower but have higher accuracy (e. g. sinf(x) and expf(x) ). The throughput of sinf(x). cosf(x). and expf(x) is much greater than that of sinf(x). cosf(x). and expf(x). The latter become even more expensive (about an order of magnitude slower) if the magnitude of the argument x needs to be reduced. Moreover, in such cases, the argument-reduction code uses local memory, which can affect performance even more because of the high latency of local memory. More details are available in the CUDA C Programming Guide . Note also that whenever sine and cosine of the same argument are computed, the sincos family of instructions should be used to optimize performance: sincosf() for single-precision fast math (see next paragraph) sincosf() for regular single-precision sincos() for double precision The - usefastmath compiler option of nvcc coerces every functionName() call to the equivalent functionName() call. This switch should be used whenever accuracy is a lesser priority than the performance. This is frequently the case with transcendental functions. Note this switch is effective only on single-precision floating point. Note: Medium Priority: Prefer faster, more specialized math functions over slower, more general ones when possible. For small integer powers (e. g. x 2 or x 3 ), explicit multiplication is almost certainly faster than the use of general exponentiation routines such as pow(). While compiler optimization improvements continually seek to narrow this gap, explicit multiplication (or the use of an equivalent purpose-built inline function or macro) can have a significant advantage. This advantage is increased when several powers of the same base are needed (e. g. where both x 2 and x 5 are calculated in close proximity), as this aids the compiler in its common sub-expression elimination (CSE) optimization. For exponentiation using base 2 or 10, use the functions exp2() or expf2() and exp10() or expf10() rather than the functions pow() or powf(). Both pow() and powf() are heavy-weight functions in terms of register pressure and instruction count due to the numerous special cases arising in general exponentiation and the difficulty of achieving good accuracy across the entire ranges of the base and the exponent. The functions exp2(). exp2f(). exp10(). and exp10f(). on the other hand, are similar to exp() and expf() in terms of performance, and can be as much as ten times faster than their pow() powf() equivalents. For exponentiation with an exponent of 13, use the cbrt() or cbrtf() function rather than the generic exponentiation functions pow() or powf(). as the former are significantly faster than the latter. Likewise, for exponentation with an exponent of -13, use rcbrt() or rcbrtf(). Replace sin(ltexprgt) with sinpi(ltexprgt). cos(ltexprgt) with cospi(ltexprgt). and sincos(ltexprgt) with sincospi(ltexprgt). This is advantageous with regard to both accuracy and performance. As a particular example, to evaluate the sine function in degrees instead of radians, use sinpi(x180.0). Similarly, the single-precision functions sinpif(). cospif(). and sincospif() should replace calls to sinf(). cosf(). and sincosf() when the function argument is of the form ltexprgt. (The performance advantage sinpi() has over sin() is due to simplified argument reduction the accuracy advantage is because sinpi() multiplies by only implicitly, effectively using an infinitely precise mathematical rather than a single - or double-precision approximation thereof.) By default, the nvcc compiler generates IEEE-compliant code for devices of compute capability 2.x, but it also provides options to generate code that somewhat less accurate but faster and that is closer to the code generated for earlier devices: - ftztrue (denormalized numbers are flushed to zero) - prec-divfalse (less precise division) - prec-sqrtfalse (less precise square root) Another, more aggressive, option is - usefastmath. which coerces every functionName() call to the equivalent functionName() call. This makes the code run faster at the cost of diminished precision and accuracy. See Math Libraries. Note: High Priority: Minimize the use of global memory. Prefer shared memory access where possible. Memory instructions include any instruction that reads from or writes to shared, local, or global memory. When accessing uncached local or global memory, there are 400 to 600 clock cycles of memory latency. As an example, the assignment operator in the following sample code has a high throughput, but, crucially, there is a latency of 400 to 600 clock cycles to read data from global memory: Much of this global memory latency can be hidden by the thread scheduler if there are sufficient independent arithmetic instructions that can be issued while waiting for the global memory access to complete. However, it is best to avoid accessing global memory whenever possible. Note: High Priority: Avoid different execution paths within the same warp. Any flow control instruction ( if. switch. do. for. while ) can significantly affect the instruction throughput by causing threads of the same warp to diverge that is, to follow different execution paths. If this happens, the different execution paths must be serialized, since all of the threads of a warp share a program counter this increases the total number of instructions executed for this warp. When all the different execution paths have completed, the threads converge back to the same execution path. To obtain best performance in cases where the control flow depends on the thread ID, the controlling condition should be written so as to minimize the number of divergent warps. This is possible because the distribution of the warps across the block is deterministic as mentioned in SIMT Architecture of the CUDA C Programming Guide. A trivial example is when the controlling condition depends only on ( threadIdx WSIZE ) where WSIZE is the warp size. In this case, no warp diverges because the controlling condition is perfectly aligned with the warps. Note: Low Priority: Make it easy for the compiler to use branch predication in lieu of loops or control statements. Sometimes, the compiler may unroll loops or optimize out if or switch statements by using branch predication instead. In these cases, no warp can ever diverge. The programmer can also control loop unrolling using For more information on this pragma, refer to the CUDA C Programming Guide. When using branch predication, none of the instructions whose execution depends on the controlling condition is skipped. Instead, each such instruction is associated with a per-thread condition code or predicate that is set to true or false according to the controlling condition. Although each of these instructions is scheduled for execution, only the instructions with a true predicate are actually executed. Instructions with a false predicate do not write results, and they also do not evaluate addresses or read operands. The compiler replaces a branch instruction with predicated instructions only if the number of instructions controlled by the branch condition is less than or equal to a certain threshold: If the compiler determines that the condition is likely to produce many divergent warps, this threshold is 7 otherwise it is 4. Note: Low Medium Priority: Use signed integers rather than unsigned integers as loop counters. In the C language standard, unsigned integer overflow semantics are well defined, whereas signed integer overflow causes undefined results. Therefore, the compiler can optimize more aggressively with signed arithmetic than it can with unsigned arithmetic. This is of particular note with loop counters: since it is common for loop counters to have values that are always positive, it may be tempting to declare the counters as unsigned. For slightly better performance, however, they should instead be declared as signed. For example, consider the following code: Here, the sub-expression stridei could overflow a 32-bit integer, so if i is declared as unsigned, the overflow semantics prevent the compiler from using some optimizations that might otherwise have applied, such as strength reduction. If instead i is declared as signed, where the overflow semantics are undefined, the compiler has more leeway to use these optimizations. Note: High Priority: Avoid the use of syncthreads() inside divergent code. Synchronizing threads inside potentially divergent code (e. g. a loop over an input array) can cause unanticipated errors. Care must be taken to ensure that all threads are converged at the point where syncthreads() is called. The following example illustrates how to do this properly for 1D blocks: In this example, the loop has been carefully written to have the same number of iterations for each thread, avoiding divergence ( imax is the number of elements rounded up to a multiple of the block size). Guards have been added inside the loop to prevent out-of-bound accesses. At the point of the syncthreads(). all threads are converged. Similar care must be taken when invoking syncthreads() from a device function called from potentially divergent code. A straightforward method of solving this issue is to call the device function from non-divergent code and pass a threadactive flag as a parameter to the device function. This threadactive flag would be used to indicate which threads should participate in the computation inside the device function, allowing all threads to participate in the syncthreads(). Having completed the GPU acceleration of one or more components of the application it is possible to compare the outcome with the original expectation. Recall that the initial assess step allowed the developer to determine an upper bound for the potential speedup attainable by accelerating given hotspots. Before tackling other hotspots to improve the total speedup, the developer should consider taking the partially parallelized implementation and carry it through to production. This is important for a number of reasons for example, it allows the user to profit from their investment as early as possible (the speedup may be partial but is still valuable), and it minimizes risk for the developer and the user by providing an evolutionary rather than revolutionary set of changes to the application. With each generation of NVIDIA processors, new features are added to the GPU that CUDA can leverage. Consequently, its important to understand the characteristics of the architecture. Programmers should be aware of two version numbers. The first is the compute capability. and the second is the version number of the CUDA Runtime and CUDA Driver APIs. The compute capability describes the features of the hardware and reflects the set of instructions supported by the device as well as other specifications, such as the maximum number of threads per block and the number of registers per multiprocessor. Higher compute capability versions are supersets of lower (that is, earlier) versions, so they are backward compatible. The compute capability of the GPU in the device can be queried programmatically as illustrated in the deviceQuery CUDA Sample. The output for that program is shown in Figure 12. This information is obtained by calling cudaGetDeviceProperties() and accessing the information in the structure it returns. Figure 12. Sample CUDA configuration data reported by deviceQuery The major and minor revision numbers of the compute capability are shown on the third and fourth lines of Figure 12. Device 0 of this system has compute capability 1.1. More details about the compute capabilities of various GPUs are in CUDA-Enabled GPUs and Compute Capabilities of the CUDA C Programming Guide. In particular, developers should note the number of multiprocessors on the device, the number of registers and the amount of memory available, and any special capabilities of the device. Certain hardware features are not described by the compute capability. For example, the ability to overlap kernel execution with asynchronous data transfers between the host and the device is available on most but not all GPUs with compute capability 1.1. In such cases, call cudaGetDeviceProperties() to determine whether the device is capable of a certain feature. For example, the asyncEngineCount field of the device property structure indicates whether overlapping kernel execution and data transfers is possible (and, if so, how many concurrent transfers are possible) likewise, the canMapHostMemory field indicates whether zero-copy data transfers can be performed. The CUDA Driver API and the CUDA Runtime are two of the programming interfaces to CUDA. Their version number enables developers to check the features associated with these APIs and decide whether an application requires a newer (later) version than the one currently installed. This is important because the CUDA Driver API is backward compatible but not forward compatible. meaning that applications, plug-ins, and libraries (including the CUDA Runtime) compiled against a particular version of the Driver API will continue to work on subsequent (later) driver releases. However, applications, plug-ins, and libraries (including the CUDA Runtime) compiled against a particular version of the Driver API may not work on earlier versions of the driver, as illustrated in Figure 13. Figure 13. Compatibility of CUDA versions When in doubt about the compute capability of the hardware that will be present at runtime, it is best to assume a compute capability of 2.0 as defined in the CUDA C Programming Guide section on Technical and Feature Specifications. To target specific versions of NVIDIA hardware and CUDA software, use the - arch. - code. and - gencode options of nvcc. Code that uses the warp shuffle operation, for example, must be compiled with - archsm30 (or higher compute capability). See Building for Maximum Compatibility for further discussion of the flags used for building code for multiple generations of CUDA-capable device simultaneously. The host runtime component of the CUDA software environment can be used only by host functions. It provides functions to handle the following: Device management Context management Memory management Code module management Execution control Texture reference management Interoperability with OpenGL and Direct3D As compared to the lower-level CUDA Driver API, the CUDA Runtime greatly eases device management by providing implicit initialization, context management, and device code module management. The CC host code generated by nvcc utilizes the CUDA Runtime, so applications that link to this code will depend on the CUDA Runtime similarly, any code that uses the cuBLAS. cuFFT. and other CUDA Toolkit libraries will also depend on the CUDA Runtime, which is used internally by these libraries. The functions that make up the CUDA Runtime API are explained in the CUDA Toolkit Reference Manual. The CUDA Runtime handles kernel loading and setting up kernel parameters and launch configuration before the kernel is launched. The implicit driver version checking, code initialization, CUDA context management, CUDA module management (cubin to function mapping), kernel configuration, and parameter passing are all performed by the CUDA Runtime. It comprises two principal parts: A C-style function interface ( cudaruntimeapi. h ). C-style convenience wrappers ( cudaruntime. h ) built on top of the C-style functions. For more information on the Runtime API, refer to CUDA C Runtime of the CUDA C Programming Guide. When deploying a CUDA application, it is often desirable to ensure that the an application will continue to function properly even if the target machine does not have a CUDA-capable GPU andor a sufficient version of the NVIDIA Driver installed. (Developers targeting a single machine with known configuration may choose to skip this section.) Detecting a CUDA-Capable GPU When an application will be deployed to target machines of arbitraryunknown configuration, the application should explicitly test for the existence of a CUDA-capable GPU in order to take appropriate action when no such device is available. The cudaGetDeviceCount() function can be used to query for the number of available devices. Like all CUDA Runtime API functions, this function will fail gracefully and return cudaErrorNoDevice to the application if there is no CUDA-capable GPU or cudaErrorInsufficientDriver if there is not an appropriate version of the NVIDIA Driver installed. If cudaGetDeviceCount() reports an error, the application should fall back to an alternative code path. A system with multiple GPUs may contain GPUs of different hardware versions and capabilities. When using multiple GPUs from the same application, it is recommended to use GPUs of the same type, rather than mixing hardware generations. The cudaChooseDevice() function can be used to select the device that most closely matches a desired set of features. Detecting Hardware and Software Configuration When an application depends on the availability of certain hardware or software capabilities to enable certain functionality, the CUDA API can be queried for details about the configuration of the available device and for the installed software versions. The cudaGetDeviceProperties() function reports various features of the available devices, including the CUDA Compute Capability of the device (see also the Compute Capabilities section of the CUDA C Programming Guide ). See CUDA Runtime and Driver API Version for details on how to query the available CUDA software API versions. All CUDA Runtime API calls return an error code of type cudaErrort the return value will be equal to cudaSuccess if no errors have occurred. (The exceptions to this are kernel launches, which return void, and cudaGetErrorString(). which returns a character string describing the cudaErrort code that was passed into it.) The CUDA Toolkit libraries ( cuBLAS. cuFFT. etc.) likewise return their own sets of error codes. Since some CUDA API calls and all kernel launches are asynchronous with respect to the host code, errors may be reported to the host asynchronously as well often this occurs the next time the host and device synchronize with each other, such as during a call to cudaMemcpy() or to cudaDeviceSynchronize(). Always check the error return values on all CUDA API functions, even for functions that are not expected to fail, as this will allow the application to detect and recover from errors as soon as possible should they occur. Applications that do not check for CUDA API errors could at times run to completion without having noticed that the data calculated by the GPU is incomplete, invalid, or uninitialized. Note: The CUDA Toolkit Samples provide several helper functions for error checking with the various CUDA APIs these helper functions are located in the samplescommoninchelpercuda. h file in the CUDA Toolkit. Each generation of CUDA-capable device has an associated compute capability version that indicates the feature set supported by the device (see CUDA Compute Capability ). One or more compute capability versions can be specified to the nvcc compiler while building a file compiling for the native compute capability for the target GPU(s) of the application is important to ensure that application kernels achieve the best possible performance and are able to use the features that are available on a given generation of GPU. When an application is built for multiple compute capabilities simultaneously (using several instances of the - gencode flag to nvcc), the binaries for the specified compute capabilities are combined into the executable, and the CUDA Driver selects the most appropriate binary at runtime according to the compute capability of the present device. If an appropriate native binary ( cubin ) is not available, but the intermediate PTX code (which targets an abstract virtual instruction set and is used for forward-compatibility) is available, then the kernel will be compiled Just In Time (JIT) (see Compiler JIT Cache Management Tools ) from the PTX to the native cubin for the device. If the PTX is also not available, then the kernel launch will fail. Alternatively, the nvcc command-line option - archsmXX can be used as a shorthand equivalent to the following more explicit - gencode command-line options described above: However, while the - archsmXX command-line option does result in inclusion of a PTX back-end target by default (due to the codecomputeXX target it implies), it can only specify a single target cubin architecture at a time, and it is not possible to use multiple - arch options on the same nvcc command line, which is why the examples above use - gencode explicitly. CUDA applications are built against the CUDA Runtime library, which handles device, memory, and kernel management. Unlike the CUDA Driver. the CUDA Runtime guarantees neither forward nor backward binary compatibility across versions. It is therefore best to redistribute the CUDA Runtime library with the application when using dynamic linking or else to statically link against the CUDA Runtime. This will ensure that the executable will be able to run even if the user does not have the same CUDA Toolkit installed that the application was built against. Note: When statically linking to the CUDA Runtime, multiple versions of the runtime can peacably coexist in the same application process simultaneously for example, if an application uses one version of the CUDA Runtime, and a plugin to that application is statically linked to a different version, that is perfectly acceptable, as long as the installed NVIDIA Driver is sufficient for both. Statically-linked CUDA Runtime The easiest option is to statically link against the CUDA Runtime. This is the default if using nvcc to link in CUDA 5.5 and later. Static linking makes the executable slightly larger, but it ensures that the correct version of runtime library functions are included in the application binary without requiring separate redistribution of the CUDA Runtime library. Dynamically-linked CUDA Runtime If static linking against the CUDA Runtime is impractical for some reason, then a dynamically-linked version of the CUDA Runtime library is also available. (This was the default and only option provided in CUDA versions 5.0 and earlier.) To use dynamic linking with the CUDA Runtime when using the nvcc from CUDA 5.5 or later to link the application, add the --cudartshared flag to the link command line otherwise the statically-linked CUDA Runtime library is used by default. After the application is dynamically linked against the CUDA Runtime, this version of the runtime library should be bundled with the application. It can be copied into the same directory as the application executable or into a subdirectory of that installation path. Other CUDA Libraries Although the CUDA Runtime provides the option of static linking, the other libraries included in the CUDA Toolkit (cuBLAS, cuFFT, etc.) are available only in dynamically-linked form. As with the dynamically-linked version of the CUDA Runtime library. these libraries should be bundled with the application executable when distributing that application. The CUDA Toolkits End-User License Agreement (EULA) allows for redistribution of many of the CUDA libraries under certain terms and conditions. This allows applications that depend on these libraries to redistribute the exact versions of the libraries against which they were built and tested, thereby avoiding any trouble for end users who might have a different version of the CUDA Toolkit (or perhaps none at all) installed on their machines. Please refer to the EULA for details. Note: This does not apply to the NVIDIA Driver the end user must still download and install an NVIDIA Driver appropriate to their GPU(s) and operating system. When redistributing the dynamically-linked versions of one or more CUDA libraries, it is important to identify the exact files that need to be redistributed. The following examples use the cuBLAS library from CUDA Toolkit 5.5 as an illustration: In a shared library on Linux, there is a string field called the SONAME that indicates the binary compatibility level of the library. The SONAME of the library against which the application was built must match the filename of the library that is redistributed with the application. For example, in the standard CUDA Toolkit installation, the files libcublas. so and libcublas. so.5.5 are both symlinks pointing to a specific build of cuBLAS, which is named like libcublas. so.5.5. x . where x is the build number (e. g. libcublas. so.5.5.17 ). However, the SONAME of this library is given as libcublas. so.5.5 : Because of this, even if - lcublas (with no version number specified) is used when linking the application, the SONAME found at link time implies that libcublas. so.5.5 is the name of the file that the dynamic loader will look for when loading the application and therefore must be the name of the file (or a symlink to the same) that is redistributed with the application. The ldd tool is useful for identifying the exact filenames of the libraries that the application expects to find at runtime as well as the path, if any, of the copy of that library that the dynamic loader would select when loading the application given the current library search path: In a shared library on Mac OS X, there is a field called the install name that indicates the expected installation path and filename the library the CUDA libraries also use this filename to indicate binary compatibility. The value of this field is propagated into an application built against the library and is used to locate the library of the correct version at runtime. For example, if the install name of the cuBLAS library is given as rpathlibcublas.5.5.dylib. then the library is version 5.5 and the copy of this library redistributed with the application must be named libcublas.5.5.dylib. even though only - lcublas (with no version number specified) is used at link time. Furthermore, this file should be installed into the rpath of the application see Where to Install Redistributed CUDA Libraries. To view a librarys install name, use the otool - L command: The binary compatibility version of the CUDA libraries on Windows is indicated as part of the filename. For example, a 64-bit application linked to cuBLAS 5.5 will look for cublas6455.dll at runtime, so this is the file that should be redistributed with that application, even though cublas. lib is the file that the application is linked against. For 32-bit applications, the file would be cublas3255.dll. To verify the exact DLL filename that the application expects to find at runtime, use the dumpbin tool from the Visual Studio command prompt: Once the correct library files are identified for redistribution, they must be configured for installation into a location where the application will be able to find them. On Windows, if the CUDA Runtime or other dynamically-linked CUDA Toolkit library is placed in the same directory as the executable, Windows will locate it automatically. On Linux and Mac, the - rpath linker option should be used to instruct the executable to search its local path for these libraries before searching the system paths: Note: It may be necessary to adjust the value of - ccbin to reflect the location of your Visual Studio installation. To specify an alternate path where the libraries will be distributed, use linker options similar to those below: For Linux and Mac, the - rpath option is used as before. For Windows, the DELAY option is used this requires that the application call SetDllDirectory() before the first call to any CUDA API function in order to specify the directory containing the CUDA DLLs. Note: For Windows 8, SetDefaultDLLDirectories() and AddDllDirectory() should be used instead of SetDllDirectory(). Please see the MSDN documentation for these routines for more information. The NVIDIA System Management Interface ( nvidia-smi ) is a command line utility that aids in the management and monitoring of NVIDIA GPU devices. This utility allows administrators to query GPU device state and, with the appropriate privileges, permits administrators to modify GPU device state. nvidia-smi is targeted at Tesla and certain Quadro GPUs, though limited support is also available on other NVIDIA GPUs. nvidia-smi ships with NVIDIA GPU display drivers on Linux, and with 64-bit Windows Server 2008 R2 and Windows 7. nvidia-smi can output queried information as XML or as human-readable plain text either to standard output or to a file. See the nvidia-smi documenation for details. Please note that new versions of nvidia-smi are not guaranteed to be backward-compatible with previous versions. ECC error counts Both correctable single-bit and detectable double-bit errors are reported. Error counts are provided for both the current boot cycle and the lifetime of the GPU. GPU utilization Current utilization rates are reported for both the compute resources of the GPU and the memory interface. Active compute process The list of active processes running on the GPU is reported, along with the corresponding process nameID and allocated GPU memory. Clocks and performance state Max and current clock rates are reported for several important clock domains, as well as the current GPU performance state ( pstate ). Temperature and fan speed The current GPU core temperature is reported, along with fan speeds for products with active cooling. Power management The current board power draw and power limits are reported for products that report these measurements. Identification Various dynamic and static information is reported, including board serial numbers, PCI device IDs, VBIOSInforom version numbers and product names. ECC mode Enable and disable ECC reporting. ECC reset Clear single-bit and double-bit ECC error counts. Compute mode Indicate whether compute processes can run on the GPU and whether they run exclusively or concurrently with other compute processes. Persistence mode Indicate whether the NVIDIA driver stays loaded when no applications are connected to the GPU. It is best to enable this option in most circumstances. GPU reset Reinitialize the GPU hardware and software state via a secondary bus reset. The NVIDIA Management Library (NVML) is a C-based interface that provides direct access to the queries and commands exposed via nvidia-smi intended as a platform for building 3rd-party system management applications. The NVML API is available on the NVIDIA developer website as part of the Tesla Deployment Kit through a single header file and is accompanied by PDF documentation, stub libraries, and sample applications see developer. nvidiatesla-deployment-kit. Each new version of NVML is backward-compatible. An additional set of Perl and Python bindings are provided for the NVML API. These bindings expose the same features as the C-based interface and also provide backwards compatibility. The Perl bindings are provided via CPAN and the Python bindings via PyPI. All of these products ( nvidia-smi. NVML, and the NVML language bindings) are updated with each new CUDA release and provide roughly the same functionality. Managing your GPU cluster will help achieve maximum GPU utilization and help you and your users extract the best possible performance. Many of the industrys most popular cluster management tools now support CUDA GPUs via NVML. For a listing of some of these tools, see developer. nvidiacluster-management. Any PTX device code loaded by an application at runtime is compiled further to binary code by the device driver. This is called just-in-time compilation ( JIT ). Just-in-time compilation increases application load time but allows applications to benefit from latest compiler improvements. It is also the only way for applications to run on devices that did not exist at the time the application was compiled. When JIT compilation of PTX device code is used, the NVIDIA driver caches the resulting binary code on disk. Some aspects of this behavior such as cache location and maximum cache size can be controlled via the use of environment variables see Just in Time Compilation of the CUDA C Programming Guide. It is possible to rearrange the collection of installed CUDA devices that will be visible to and enumerated by a CUDA application prior to the start of that application by way of the CUDAVISIBLEDEVICES environment variable. Devices to be made visible to the application should be included as a comma-separated list in terms of the system-wide list of enumerable devices. For example, to use only devices 0 and 2 from the system-wide list of devices, set CUDAVISIBLEDEVICES0,2 before launching the application. The application will then enumerate these devices as device 0 and device 1, respectively. This appendix contains a summary of the recommendations for optimization that are explained in this document. Performance optimization revolves around three basic strategies: Maximizing parallel execution Optimizing memory usage to achieve maximum memory bandwidth Optimizing instruction usage to achieve maximum instruction throughput Maximizing parallel execution starts with structuring the algorithm in a way that exposes as much data parallelism as possible. Once the parallelism of the algorithm has been exposed, it needs to be mapped to the hardware as efficiently as possible. This is done by carefully choosing the execution configuration of each kernel launch. The application should also maximize parallel execution at a higher level by explicitly exposing concurrent execution on the device through streams, as well as maximizing concurrent execution between the host and the device. Optimizing memory usage starts with minimizing data transfers between the host and the device because those transfers have much lower bandwidth than internal device data transfers. Kernel access to global memory also should be minimized by maximizing the use of shared memory on the device. Sometimes, the best optimization might even be to avoid any data transfer in the first place by simply recomputing the data whenever it is needed. The effective bandwidth can vary by an order of magnitude depending on the access pattern for each type of memory. The next step in optimizing memory usage is therefore to organize memory accesses according to the optimal memory access patterns. This optimization is especially important for global memory accesses, because latency of access costs hundreds of clock cycles. Shared memory accesses, in counterpoint, are usually worth optimizing only when there exists a high degree of bank conflicts. As for optimizing instruction usage, the use of arithmetic instructions that have low throughput should be avoided. This suggests trading precision for speed when it does not affect the end result, such as using intrinsics instead of regular functions or single precision instead of double precision. Finally, particular attention must be paid to control flow instructions due to the SIMT (single instruction multiple thread) nature of the device. The NVIDIA nvcc compiler driver converts. cu files into C for the host system and CUDA assembly or binary instructions for the device. It supports a number of command-line parameters, of which the following are especially useful for optimization and related best practices: - maxrregcountN specifies the maximum number of registers kernels can use at a per-file level. See Register Pressure. (See also the launchbounds qualifier discussed in Execution Configuration of the CUDA C Programming Guide to control the number of registers used on a per-kernel basis.) --ptxas-options-v or - Xptxas-v lists per-kernel register, shared, and constant memory usage. - ftztrue (denormalized numbers are flushed to zero) - prec-divfalse (less precise division) - prec-sqrtfalse (less precise square root) - usefastmath compiler option of nvcc coerces every functionName() call to the equivalent functionName() call. This makes the code run faster at the cost of diminished precision and accuracy. See Math Libraries. ALL NVIDIA DESIGN SPECIFICATIONS, REFERENCE BOARDS, FILES, DRAWINGS, DIAGNOSTICS, LISTS, AND OTHER DOCUMENTS (TOGETHER AND SEPARATELY, MATERIALS) ARE BEING PROVIDED AS IS. NVIDIA MAKES NO WARRANTIES, EXPRESSED, IMPLIED, STATUTORY, OR OTHERWISE WITH RESPECT TO THE MATERIALS, AND EXPRESSLY DISCLAIMS ALL IMPLIED WARRANTIES OF NONINFRINGEMENT, MERCHANTABILITY, AND FITNESS FOR A PARTICULAR PURPOSE. Information furnished is believed to be accurate and reliable. However, NVIDIA Corporation assumes no responsibility for the consequences of use of such information or for any infringement of patents or other rights of third parties that may result from its use. No license is granted by implication of otherwise under any patent rights of NVIDIA Corporation. Specifications mentioned in this publication are subject to change without notice. This publication supersedes and replaces all other information previously supplied. NVIDIA Corporation products are not authorized as critical components in life support devices or systems without express written approval of NVIDIA Corporation. Trademarks NVIDIA and the NVIDIA logo are trademarks or registered trademarks of NVIDIA Corporation in the U. S. and other countries. Other company and product names may be trademarks of the respective companies with which they are associated. The surprise appearance of a tractor-trailer rig on stage at AWS ReInvent in Las Vegas illustrated a new way to move a petabyte of data into AWS data centers. Por Charles Babcock Editor em grande, Cloud, 152017 As conseqüências políticas e técnicas do relatório conjunto do DHS-FBI Grizzly Steppe sobre o papel das Russias nos recentes hacks relacionados a eleições causam mais caos do que o encerramento. Por Kelly Jackson Higgins Editor Executivo em Dark Reading, 142017 A conformidade com a privacidade é agora uma missão crítica. Os fornecedores de terceiros que não atendem aos mandatos de proteção de dados serão excluídos de fazer negócios em lucrativos mercados verticais. Por John Moynihan Presidente, Minuteman Governance, 142017 O registrador de domínio e provedor de hospedagem web GoDaddy gera muitos dados, e queria ajudar seus usuários internos a obter melhores informações sem requerer intervenção da equipe técnica. Por Jessica Davis Editor sênior, Enterprise Apps, 142017 O setor de saúde tem atuado como indústrias como bancos e varejo quando se trata de adoção de análise de dados, mas um número crescente de histórias de sucesso na área de saúde fornece prova de conceito para outras organizações para obter borda. Por Jeremy Achin, CEO da DataRobot. 142017 A Theres foi uma mudança positiva no negócio de informática empresarial de anúncios de produtos mais recentes, verdadeiramente focados em como a tecnologia pode ajudar uma empresa a fazer negócios. Por James M. Connolly Editor Executivo Executivo, InformationWeekEditor em Chefe, 132017 A TSO Logic examinou dados de 10.000 servidores de clientes e descobriu quanto as atualizações podem economizar no tamanho do direito de contagem de servidores economiza recursos de VM. Por Charles Babcock Editor em grande, Cloud, 132017
No comments:
Post a Comment