Como parte da coluna "Faça uma pergunta a um especialista da Intel", solicitamos ao especialista em Intel Konstantin Vladimirov que respondesse perguntas relacionadas à programação heterogênea, ao
kit de ferramentas oneAPI e a coisas interessantes relacionadas. O resultado superou todas as nossas expectativas. Konstantin não economizou tempo e deu respostas detalhadas e substanciadas, sem medo de ser polêmico. De fato, recebemos uma pequena palestra sobre programação entre arquiteturas em todas as suas formas: descarregar nuances, otimizações, padrões e assim por diante.
Transferimos o microfone para o especialista. Bem, os comentários são dados ao público.
Pergunta
Soarex16Quão trabalhosa será a transição do OpenCL para o oneAPI e quais benefícios podem ser obtidos com isso?
A resposta Mudar para o DPC ++ pode ser complicado, mas, na minha opinião, vale a pena. Existem dois estágios principais.
Primeiramente, essa é uma transição da sua linguagem de programação heterogênea (OpenCL, computação Vulkan), que, provavelmente, é baseada na API. Aqui você tem uma vantagem no fato de que você já conhece a área de assunto e a dificuldade está em mudar o pensamento do controle direto via API para construções de linguagem um pouco mais implícitas.
Em segundo lugar, essa é uma transição do seu idioma host. Se você descarregou toda a sua vida do C puro, o limite de entrada é igual ao limite para alternar de C para C ++, que é bastante alto.
Por que tentar?
Em primeiro lugar, o DPC ++ faz um ótimo trabalho para um programador. Você esquecerá muito rapidamente, como um pesadelo, todas essas chamadas explícitas para clXXXYYY, e o que o sexto argumento significa, e se você esqueceu o código de retorno. Muitos wrappers orientados a objetos não escondem a rotina, mas geralmente com o custo de alternar da API OpenCL padrão para a API wrapper não tão padrão (eu também vi essas bicicletas). No caso do DPC ++, você simplesmente escreve o SYCL padrão com extensões Intel (que em breve também se tornará o SYCL padrão).
Em segundo lugar, o DPC ++ fornece compilação conjunta, ou seja, você pode ter certeza dos tipos e não terá problemas nas bordas da API com dimensões, preenchimento, alinhamento. Você escreve o código do kernel e do host em um arquivo e esse é o mesmo código. Usando o USM, você também pode trabalhar com estruturas de dados complexas com muito mais facilidade.
Em terceiro lugar, o DPC ++ é C ++ real, ou seja, permite programação generalizada. Por exemplo, o kernel mais simples para adicionar dois vetores:
auto kern = [A, B, C](cl::sycl::id<1> wiID) { C[wiID] = A[wiID] + B[wiID];
A mesma coisa no OpenCL:
_kernel void vector_add(__global int *A, __global int *B, __global int *C) { int i = get_global_id(0); C[i] = A[i] + B[i]; }
Você vê, fui forçado a apontar para um tipo OpenCL int. Se eu precisar de um float, precisarei escrever outro kernel ou usar um pré-processador ou geração de código externo. Obter quase todos os recursos do C ++ à sua disposição pode ser um pouco assustador se você não tiver experiência com o C ++. Mas isso é comum quando se trata de uma grande mudança tecnológica.
E todos os benefícios não se limitam a isso. Mencionarei outra coisa nas respostas a seguir.
Então, eu teria baixado o compilador em seu lugar e tentado, pois não é difícil fazer isso com o pacote
OneAPI .
Question
JusterO OpenVINO e o oneAPI estarão de alguma forma relacionados?
A resposta A distribuição OpenVINO agora faz parte da distribuição OneAPI. Aprender e usar redes neurais são tarefas computacionalmente difíceis que se beneficiam muito da programação heterogênea. Acredito que, mais cedo ou mais tarde, todos os componentes OneAPI tornarão possível o uso de todos os recursos de computação disponíveis: aceleradores gráficos e aceleradores especiais como Nervana e FPGA. E tudo isso sem sair do paradigma da linguagem e do sistema de tipos do seu programa C ++.
Perguntas do correio
Estou tentando entender como o acelerador de hardware de IA ficará em três anos. Por favor, ajude com isso. Existe uma empresa interessante Graphcore e sua IPU - este dispositivo não é menos eficiente que o FPGA, mas é muito mais fácil programar - Python com suporte para TensorFlow e outras estruturas. Acontece que, se as promessas da Graphcore forem cumpridas, não haverá necessidade de FPGAs no mercado de aprendizado de máquina. O Python é muito mais conveniente para os cientistas de dados do que o C ++.
Você concorda que o FPGA não é adequado para o mercado de aprendizado de máquina comparado às soluções programáveis em Python? Se o mercado de ML é perdido, que outros aplicativos FPGA comuns você vê?
Em quais aplicativos você vê a necessidade inevitável de programação heterogênea, onde você não pode conviver com ferramentas mais convenientes, como Python?
A resposta Olhei brevemente para que tipo de IPU. Mais um pedaço de ferro no qual todos descarregarão. Esses caras competem com a GPU e com aceleradores especiais, e não com o FPGA.
Nas tarefas para as quais uma peça de hardware especializada é afiada, ela sempre vence o FPGA, por exemplo, a renderização de vídeo é melhor em uma placa de vídeo etc. Mas no mundo (inclusive no mundo da ML) existem muitas tarefas para as quais nada de especial foi inventado ou lançado, e aqui o FPGA sempre será indispensável. Por exemplo, porque existe uma questão de preço e, para ser barato, uma peça de hardware especializada deve ser enorme.
Suponha agora que a IPU especificada é realmente legal. Isso não cancelará a programação heterogênea; pelo contrário, a presença de um acelerador tão excelente o estimulará. E também dará um avanço gigantesco no OneAPI e no DPC ++, porque mais cedo ou mais tarde alguém dirá "Eu quero usar tanto a sua IPU quanto a minha GPU em um programa". Bem cedo, porque a programação heterogênea é sobre isso. Seu significado é a transferência de uma tarefa adequada para um dispositivo adequado. Uma tarefa pode vir de qualquer lugar. E esse dispositivo pode ser qualquer coisa, pode até ser o mesmo dispositivo em que o programa está sendo executado. Por exemplo, se você descarregar o kernel escrito no ISPC e utilizar os recursos vetoriais do Xeon ao máximo, poderá descarregá-lo você mesmo e ainda obter um ganho significativo. O principal critério aqui é desempenho. Bem, nunca haverá muita produtividade neste mundo. Mesmo com os melhores aceleradores do mundo.
Quanto ao Python e sua conveniência ... Devo admitir imediatamente que não gosto de linguagens dinamicamente digitadas: elas são lentas e, em vez de um erro de compilação normal, você deve esperar duas horas antes de cair no tempo de execução devido ao tipo errado. Mas não vejo como é ruim fazer as mesmas descargas no Python. A propósito, o OneAPI já inclui o Intel Distribution for Python, o que é extremamente conveniente para várias revisões.
Ou seja, no mundo dos sonhos dos amantes de Python, você escreve um programa e o transfere para todos os aceleradores que pode encontrar usando o OneAPI, e não para um monte de bibliotecas específicas de fornecedores. Outra coisa é que, com essa abordagem, você perde a digitação de ponta a ponta e retorna ao mundo extremamente precário da programação baseada em API. Talvez o desenvolvimento do DPC ++ incentive a comunidade a usar mais ativamente ferramentas mais apropriadas, como C ++.
Pergunta do correio
Desempenho versus OpenCL. Deve haver impostos sobre o luxo - ou seja, custos indiretos. Existem medidas?
A resposta Na Internet, você pode encontrar muitas medições com diversos resultados, dependendo do compilador, da tarefa e da qualidade da implementação. Como uma pesquisa pessoal, avaliei tarefas simples (SGEMM, DGEMM) no meu laptop (gráficos integrados do Skylake) e vi que, até agora, há algum problema (em porcentagem). Mas parece-me que isso é uma conseqüência do fato de que tudo isso é beta até agora.
Em teoria, o resultado deve ser aceleração, não desaceleração, ou seja, em princípio, todo esse luxo deve ter um valor negativo. É tudo sobre o compilador. Quando seu programa consiste em uma única fonte e é processado como um único programa, o compilador obtém oportunidades fantásticas e incríveis de otimização: definindo código comum, invertendo loops, reorganizando seções de código e tudo o mais que o compilador simplesmente não pode fazer na abordagem baseada em API, mas mais cedo ou mais tarde, ela definitivamente aprenderá com um único modelo de fonte.
Além disso, o DPC ++ terá um custo negativo em termos de tempo de desenvolvimento. Um exemplo simples são os acessadores SYCL, que o compilador já está usando para organizar eventos e gerenciar filas assíncronas.
deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto C = bufferC.template get_access<sycl_write>(cgh); .... deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto D = bufferD.template get_access<sycl_write>(cgh);
Aqui, o compilador vê que os dois pacotes apenas lêem A e B e escrevem buffers independentes C e D; como resultado, ele vê a capacidade de enviá-los em paralelo, se houver tamanhos globais suficientes.
Obviamente, um programa OpenCL gravado de forma pedestre também pode fazê-lo, mas o tempo de desenvolvimento gasto com um kernel não trivial não será comparável.
Pergunta do correio
Todas as maneiras de otimizar aplicativos OpenCL para DPC ++ são relevantes? O que há de novo a ser adicionado a eles?
A resposta Eu diria que a maior parte da sutil otimização manual que está sendo feita pelos gravadores do kernel pode e deve ser feita pelo compilador. Da mesma forma, por exemplo, considero uma prática prejudicial instalar manualmente um assembler em linha em programas C ++, porque mesmo que ofereça benefícios táticos, ele interfere nas otimizações e atua como um fator negativo no desenvolvimento e na transferência de um produto. Bem, o OpenCL agora também é montador.
Quanto à resposta mais detalhada, tenho medo do abismo aqui. Por exemplo, existe um documento Intel bem conhecido "Guia do desenvolvedor OpenCL para gráficos de processador Intel". E há uma
seção sobre como tentar, para não colocar onde o excesso de sincronização.
Portanto, do meu ponto de vista, essa é uma tarefa não humana em princípio. As pessoas são extremamente pobres em raciocinar sobre a sincronização multithread e tendem a esculpir a sincronização de maneira conservadora ou incorreta, ou ambas ao mesmo tempo - coloquei vírgulas como essa (
mas foi corrigida - nota editorial ).
Por outro lado, no DPC ++, em vez de escrever código com barreiras explícitas, assim:
for (t = 0; t < numTiles; t++) { const int tiledRow = TS * t + row; const int tiledCol = TS * t + col; Asub[col][row] = A[globalRow * AY + tiledCol]; Bsub[col][row] = B[tiledRow * BY + globalCol];
Você provavelmente escreverá uma iteração explícita de
parallel_for_work_group , dentro da qual
group.parallel_for_work_item cgh.parallel_for_work_group<class mxm_kernel>( cl::sycl::range<2>{BIG_AX / TS, BIG_BY / TS}, cl::sycl::range<2>{TS, TS}, [=](cl::sycl::group<2> group) {
Como resultado, você não precisa definir a sincronização manualmente, e toda a seção pode ser descartada.
E assim você pode andar em todas as seções. Algo vai sobreviver, algo vai sair. Prevejo o surgimento de um novo documento "Otimização para DPC ++", mas o tempo deve passar, pois todas as técnicas realmente úteis são desenvolvidas apenas mais tarde e com sangue
Pergunta do correio
Há uma limitação no OpenCL - você não pode usar "dados distantes" no kernel, ou seja, por exemplo, implementar um "filtro amplo" que usa dados de entrada de um grande grupo de pixels maior que o grupo de trabalho OpenCL em um cálculo. O que o DPC ++ oferece a esse respeito?
A resposta Bem, isso é impossível. Obviamente, não escrevo kernels em particular ... Mas é absolutamente certo que você pode usar toda a memória global como ela é, basta garantir que trabalhe com operações atômicas (ou sincronize externamente os kernels hierárquicos). E você também pode conectar o System SVM (bem, ou o USM no DPC ++).
Infelizmente, tudo isso é extremamente ineficiente, e eu não gosto de todos esses truques. Além disso, eles são difíceis de otimizar pelo compilador.
E assim, se falamos de soluções diretas e eficazes, é claro que não há mágica no DPC ++. No final, seu programa ainda está dividido em partes: o código do host e o código do dispositivo, e todas as restrições do dispositivo afetam o código do dispositivo. O tamanho máximo do grupo de trabalho é o paralelismo real de que seu hardware é capaz. Tudo o que está em cima disso são apenas maneiras de sair, afetando drasticamente o desempenho. É por isso que o DPC ++ oferece uma oportunidade para fazer isso:
device.get_info <sycl :: info :: device :: max_work_group_size> () e depois decide como viver com o número resultante.
Seria tentador, é claro, criar um modelo no DPC ++, quando o programador trabalha como você gosta com loops de qualquer tamanho, e o compilador examina o que fazer a seguir, mas seria mortalmente errado, porque ocultaria constantes e até assintóticos de complexidade adicional computação aparecendo do nada. Por outro motivo, Alexandrescu escreveu que "encapsular a complexidade deve ser considerado um crime", e isso também se aplica.
Às vezes, revisar o próprio algoritmo ajuda. Aqui, o DPC ++ facilita as coisas porque é mais fácil refatorar um código mais estruturado. Mas isso é tão consolador.
Pergunta do correio
O DPC ++ é baseado em SYCL. Mas e se você for mais fundo, quais são as diferenças do OpenCL na implementação do backend, se houver. Por exemplo, o mecanismo de distribuição entre dispositivos heterogêneos é o mesmo que o OpenCL?
A resposta Se você se esconder, esse é o OpenCL. Todas as vantagens e vantagens do SYCL são as vantagens e os pontos fortes do idioma, ou seja, o frontend. Do front-end, vem o bom e velho SPIRV, que vai para o back-end e lá é otimizado (geralmente já em tempo de execução, ou seja, é JIT) já para uma placa de vídeo específica da mesma maneira que o OpenCL seria otimizado para ela.
Outra coisa é que o mecanismo de distribuição de trabalho entre dispositivos heterogêneos é apenas mais front-end que back-end, porque é o código do host que decide o que enviar e para onde. E o código do host é obtido no DPC ++. Eu já mostrei um exemplo um pouco mais alto de como o compilador pode, com base nos acessadores, tomar uma decisão sobre pacotes paralelos. E esta é apenas a ponta do iceberg.
Pergunta do correio
Bibliotecas Sim, não estamos falando de CUDA. Mas sabemos que, para desenvolvedores de CUDA, existem bibliotecas muito úteis que funcionam com alto desempenho na GPU. O OneAPI também contém algumas bibliotecas, mas, por exemplo, IPP - não há nada útil de arquivamento para trabalhar com imagens no oneAPI / OpenCL. Haverá algo e como, neste caso, mudar de CUDA para oneAPI?
A resposta A transição da CUDA para um único padrão aberto será difícil, mas inevitável. Obviamente, a CUDA agora tem uma infraestrutura mais madura. Mas as características de seu licenciamento são uma desvantagem de bloqueio, porque cada vez mais players aparecem no mercado de sistemas heterogêneos, cartões e aceleradores cada vez mais interessantes de diferentes fabricantes.
A diversidade de APIs existentes dificulta o uso desse mundo de possibilidades para programadores com experiência na CPU clássica. O que leva ao OneAPI ou algo parecido. Aqui a mágica não está na inovação da Intel nos gráficos, mas no fato de que a Intel abre a porta do DPC ++ para todos. Nós nem possuímos o padrão SYCL, ele pertence ao grupo Khronos e todas as extensões Intel são extensões no Khronos onde qualquer pessoa pode se comprometer (e há representantes de todos os principais players do país). E isso significa que (as bibliotecas) e a comunidade aparecerão (já estão aparecendo) e várias vagas nessa direção.
E, é claro, o IPP será reescrito para novas realidades. Não tenho nada a ver com IPP, mas o uso do DPC ++ é senso comum, e pessoas sãs estão sentadas lá.
Mas o mais importante é que agora é o momento da história em que você pode escrever sua própria biblioteca, que ultrapassará o IPP e que o mundo inteiro usará. Porque os padrões abertos sempre vencem.
Pergunta do correio
Se compararmos o lançamento de algoritmos de redes neurais de treinamento e inferência no Nervana e FPGA - quais são as diferenças na programação e a eficiência resultante?
A resposta Não sei nada sobre detalhes de programação de FPGA, escrevo compiladores. Mas eu tenho uma contra-pergunta. E como vamos comparar? Em benchmarks padrão, é antidesportivo, Nervana lambeu embaixo deles. Mas caso você tenha algo interessante, o FPGA desamarra as suas mãos, e colocar isso em Nervana pode ser longo, caro, só isso.
Acontece que a pergunta em si é, por assim dizer, da série "quem é mais forte que um elefante ou uma baleia". Mas essa não é uma pergunta real. A verdadeira questão é: como aproveitar um elefante e uma baleia em um carrinho? Bem, ou pelo menos distribua, digamos, um elefante para puxá-lo por terra e uma baleia por mar.
No caso do OneAPI, você terá o mesmo programa, em geral, em C ++ padrão. E você pode escrevê-lo e executá-lo com descarregamento para frente e para trás. Essa será a mesma tarefa que lhe interessa, na qual você mesmo pode medir e otimizar o desempenho. Um padrão único e uma interface única para dispositivos heterogêneos serão um passo para comparar maçãs com maçãs em tais assuntos.
Por exemplo: “o que é melhor para% da minha tarefa% do ponto de vista da facilidade de programação e eficiência - coloque essa parte no FPGA, deixe essa no Nervana ou divida essa parte em duas e reescreva essa parte para a GPU?”
E a história toda com a OneAPI - é apenas para você dizer: "por que pensar nisso por um longo tempo, vou tentar agora rapidamente, é SIMPLES".
Ainda não, não é fácil. Mas haverá.
Posfácio do especialistaObrigado a todos por suas perguntas. É possível e até provável que eu estivesse errado, impreciso e cometi erros. Acontece que, na Internet, constantemente alguém está errado.
Espero poder interessar alguém em programação heterogênea e DPC ++. Quero recomendar a todos o site
sycl.tech , onde
existem muitos relatórios, inclusive de especialistas de renome mundial (é necessário inglês)
Bom para todos!
PS do editor. Desta vez, por decisão unânime do conselho editorial, foi decidido atribuir o prêmio pela melhor pergunta ... ao autor das respostas. Eu acho que você concorda que isso é justo.