Geekly articles weekly

Exemplo de programação do acelerador FPGA



Há pouco tempo, falamos sobre o novo serviço Selectel - computação em nuvem de alto desempenho em aceleradores FPGA . Em um novo artigo sobre este assunto, consideramos um exemplo de programação FPGA para a construção de um conjunto de Mandelbrot, um conhecido algoritmo matemático para visualização de imagens fractais. O artigo utilizou material do site Euler Project .



Em vez do prefácio


Primeiro, alguns termos. Um sistema de computador com um acelerador FPGA - como regra, este é um adaptador PCIe com um chip FPGA como parte do servidor x64. O acelerador assume uma tarefa separada de uso intensivo de recursos, na qual a computação paralela pode estar envolvida e executa muitas ordens de magnitude mais rapidamente que o processador x64, descarregando-o e aumentando o desempenho de todo o sistema de computação. Por exemplo, um ciclo de cálculo com 100 mil repetições pode ser realizado em um FPGA em apenas uma passagem, em vez de executar sequencialmente 100 mil vezes em um processador x64 clássico. Elementos lógicos, recursos de hardware, links de comunicação, chips FPGA são programados pelo usuário diretamente para a tarefa em si, o que permite implementar a tarefa como uma implementação de um algoritmo em silício - algoritmo em silício e, assim, alcançar alto desempenho e com um consumo de energia muito modesto.

Hoje, o limite para a entrada na tecnologia FPGA é bastante acessível, mesmo para as startups - um servidor com um acelerador FPGA e todo o software necessário (SDK) podem ser alugados na nuvem Selectel por um preço razoável (o chamado "cloud FPGA"), e o suporte ao padrão Open CL no FPGA leva a que um programador que sabe trabalhar com C é capaz de preparar e executar um programa no FPGA.

Olhando para o futuro: experimente o FPGA no trabalho


O exemplo de programação descrito abaixo para a construção de um conjunto Mandelbrot já foi implementado em um servidor de teste no laboratório Selectel , onde qualquer pessoa pode avaliar seu desempenho (o registro será necessário).


O projeto é fornecido em código e preparado para compilação. Selectel oferece acesso remoto a um servidor com o acelerador Intel Arria 10 FPGA. No lado do servidor, são implementadas as ferramentas SDK e BSP para desenvolvimento, depuração e compilação do código OpenCL, Visual Studio para preparar aplicativos host (aplicativos de controle para o processador central do servidor).
Observe que o exemplo em si não tem nenhum valor aplicado; foi escolhido por razões de demonstração dos métodos de aceleração usando os princípios do paralelismo. Com este exemplo, o leitor se familiariza com a rota de design de um aplicativo em um sistema de computação heterogêneo com FPGA - posteriormente, essa rota pode ser usada para desenvolver seus próprios aplicativos com computação paralela.
ATUALIZAÇÃO : Na primavera de 2018, a Intel lançou o processador híbrido de alto desempenho Xeon Gold 6138P com um chip FPGA Arria 10 integrado. Até o final de 2018, espera-se que os processadores seriais desse tipo estejam disponíveis para os clientes por meio de parceiros da Intel. Na Selectel, esperamos ansiosamente por esse chip e esperamos ser os primeiros na Rússia a oferecer aos nossos clientes a oportunidade de testar este novo produto exclusivo.

Sobre o padrão OpenCL para programação FPGA


O padrão OpenCL foi desenvolvido pelo Khronos Group, os principais fabricantes de chips e software do mundo, incluindo Intel, AMD, Apple, ARM, Nvidia, Sony Computer Entertainment, etc. Ele foi projetado para escrever aplicativos que usam computação paralela em vários tipos de processadores, incluindo FPGA. O padrão OpenCL inclui a linguagem de programação C baseada na versão da linguagem C99 (a versão mais recente da C99 é ISO / IEC 9899: 1999 / Cor 3: 2007 de 15-11-2007) e um ambiente de programação de aplicativos.

A popularidade do uso do OpenCL para computação de alto desempenho baseia-se no fato de ser um padrão aberto e seu uso não requer uma licença. Além disso, o OpenCL não limita a gama de dispositivos suportados a nenhuma marca em particular, permitindo o uso de hardware de diferentes fabricantes na mesma plataforma de software.

Além disso, sobre o OpenCL: Introdução ao OpenCL no Habr .


Um pouco de história - a rota de design do FPGA que existia antes do padrão OpenCL era extremamente específica e demorada, enquanto em termos de complexidade ultrapassava o design de microcircuitos personalizados (ASIC, circuito integrado específico da aplicação, "circuito integrado para fins especiais"). Foi necessário um entendimento completo da estrutura de hardware do FPGA, cuja configuração teve que ser realizada em uma linguagem de descrição de hardware (HDL) de baixo nível. A posse dessa rota de design e verificação foi e continua sendo uma arte que, devido à extrema complexidade, está disponível para um círculo limitado de desenvolvedores.

O advento do kit de ferramentas de suporte OpenCL da Intel para FPGAs abordou parcialmente a questão da acessibilidade da programação FPGA para desenvolvedores de software. O programador seleciona independentemente a parte do algoritmo que é adequada para processamento paralelo e a descreve em C; em seguida, o compilador Intel OpenCL para FPGA cria um arquivo de configuração binária para executar esse fragmento do algoritmo no acelerador.
Usando o ambiente normal do Visual Studio ou o compilador gcc padrão, um aplicativo host é preparado (um aplicativo do tipo .exe, executado no processador x64 principal), enquanto todas as bibliotecas de suporte necessárias estão incluídas no SDK. Quando o aplicativo host é iniciado, o firmware do FPGA é carregado, os dados são carregados no núcleo do chip e o processamento começa de acordo com o algoritmo concebido.

O FPGA (FPGA) é uma estrutura de hardware paralelo em massa reprogramável pelo usuário com milhões de elementos lógicos, milhares de blocos de sinal DSP e dezenas de megabytes de cache para cálculos a bordo, sem acessar os principais módulos de memória do servidor. Interfaces de E / S rápidas (10GE, 40GE, 100GE, PCIe Gen 3, etc.) permitem que você troque dados com o processador principal do servidor.


O padrão OpenCL é um ambiente para a execução de software heterogêneo. O ambiente consiste em duas partes separadas:

  1. Software host - um aplicativo em execução no processador central principal do servidor, escrito em C / C ++ e usando o conjunto de funções da API OpenCL. O servidor host organiza todo o processo de computação, fornecendo a fonte e recebendo dados de saída e interage com todos os sistemas de servidor com o acelerador FPGA.
  2. Software Accelerator - um programa escrito na linguagem C do OpenCL (linguagem C com várias restrições), compilado para ser executado no chip FPGA.

Um servidor típico para computação paralela é um computador baseado em x64 (para executar aplicativos host), que inclui um acelerador FPGA de hardware, geralmente conectado via barramento PCI-Express. A propósito, apenas esse sistema é apresentado no laboratório Selectel.

A sequência de programação e compilação para o acelerador FPGA consiste em dois estágios. O código do aplicativo host é compilado por um compilador padrão (Visual C ++, GCC) para obter um arquivo executável no sistema operacional do servidor (por exemplo, * .exe). O código fonte do acelerador FPGA (kernel, kernel) é preparado pelo compilador AOC como parte do SDK, com o recebimento de um arquivo binário (* .aocx). Este arquivo é apenas para programação do acelerador.

Arquitetura do ambiente de compilação de software OpenCL
Fig. Arquitetura do ambiente de compilação de software OpenCL

Considere um código de exemplo para calcular um vetor grande de duas maneiras
( PS Não atire no pianista - daqui em diante o código do site do Projeto Euler é usado ):

void inc (float *a, float c, int N) { for (int i = 0; i<N; i++) a[i] = a[i] + c; } void main() { ... inc(a,c,N); ... } 

 _kernel void inc (_global float *a, float c) { int i = get_global_id(0); a[i] = a[i] + c; } void main() { ... clEnqueueNDRangeKernel(...,&N,...) ... }

O código no início é um exemplo de como uma implementação de thread único em C pode parecer usando o método de cálculo seqüencial de elementos escalares.

A segunda versão do código é uma possível implementação do algoritmo no OpenCL na forma de uma função calculada em um acelerador FPGA. Não há loop, e o cálculo ocorre em uma iteração do loop. O cálculo de uma matriz de vetores ocorre como a execução de N cópias desta função. Cada cópia possui seu próprio índice, substituído no iterador em um loop, e o número de tentativas é definido no host quando o código é executado. A ação do iterador é fornecida pela função get_global_id (), que trabalha com um índice dentro de 0 ≤ index <N.

Vá direto ao ponto: construindo um fractal


O conjunto de Mandelbrot é uma matriz de pontos “c” no plano complexo para o qual a relação de recorrência Zn + 1 = Zn² + c para Z0 = 0 define uma sequência limitada.

Definimos Zn = Zn + IYn, e também c = p + iq.
Para cada ponto, a seguinte sequência é calculada:

Xn + 1 = Xn² + Yn² + p
Yn + 1 = 2XnYn + q


O cálculo da pertença de um ponto ao conjunto em cada iteração é realizado como a equação
Xn² + Yn² <4.

Para exibir o conjunto de Mandelbrot na tela, definimos uma regra:

  1. Se a desigualdade se mantiver em qualquer iteração, o ponto entra no conjunto e será mostrado em preto.
  2. Se a desigualdade não se mantiver, começando com um certo valor de iteração n = N, a cor é determinada pelo número de iterações N.

O processo de cálculo no host será o seguinte:

  • O cálculo do número de iterações para cada ponto dentro da janela de pixel é atribuído à função mandel_pixel ().
  • A enumeração seqüencial dos pontos da imagem será fornecida pela função softwareCalculateFrame (). Os parâmetros especificam o intervalo real dos pontos calculados, a etapa real do algoritmo e um ponteiro para o buffer de cores do tamanho da imagem (theWidth * theHeight).
  • A cor do ponto é ajustada pelaSoftColorTable.

Vamos para o código:

 inline unsigned int mandel_pixel( double x0, double y0, unsigned int maxIterations ) { // variables for the calculation double x = 0.0; double y = 0.0; double xSqr = 0.0; double ySqr = 0.0; unsigned int iterations = 0; // perform up to the maximum number of iterations to solve // the current point in the image while ( xSqr + ySqr < 4.0 &&iterations < maxIterations ) { // perform the current iteration xSqr = x*x; ySqr = y*y; y = 2*x*y + y0; x = xSqr - ySqr + x0; // increment iteration count iterations++; } // return the iteration count return iterations; } 

 int softwareCalculateFrame( double aStartX, double aStartY, double aScale, unsigned int* aFrameBuffer ) { // temporary pointer and index variables unsigned int * fb_ptr = aFrameBuffer; unsigned int j, k, pixel; // window position variables double x = aStartX; double y = aStartY; double cur_x, cur_y; double cur_step_size = aScale; // for each pixel in the y dimension window for ( j = 0, cur_y = y; j < theHeight; j++, cur_y -= cur_step_size ) { // for each pixel in the x dimension of the window for ( cur_x = x, k = 0; k< theWidth; k++, cur_x += cur_step_size ) { // set the value of the pixel in the window pixel = mandel_pixel(cur_x, cur_y, theSoftColorTableSize); if ( pixel == theSoftColorTableSize ) *fb_ptr++ = 0x0; else *fb_ptr++ = theSoftColorTable[pixel]; } } return 0; }

Cada pixel é calculado independentemente do outro e, portanto, esse processo pode ser paralelo. Ao implementar o algoritmo para o acelerador FPGA, uma instrução SIMD é criada para calcular o número de cada pixel de iteração (determinando o código de cores da paleta). A implementação de dois loops aninhados no buffer de imagem é enquadrada pelo OpenCL executando a operação (theWidth * theHeight).

As instâncias do kernel na lista abaixo são chamadas de item de trabalho e o conjunto de todas as instâncias é chamado de espaço de índice. Os recursos da função de hardware incluem o seguinte:

  • Uma declaração de função começa com a palavra-chave __kernel.
  • Tipo de função de hardware - o tipo do valor de retorno é sempre nulo.
  • O retorno de valores é feito através de buffers passados ​​como parâmetros.
    • Os três primeiros parâmetros definem a grade do material, cujos nós correspondem aos pixels da imagem de saída.
    • O quarto parâmetro limita o número de iterações, impedindo o loop de pontos pertencentes ao conjunto Mandelbrot.
    • O quinto parâmetro é um ponteiro para o buffer de cores de saída.
    • A palavra-chave global indica o tipo de memória através da qual o buffer será transmitido: esta é a memória DDR geral (QDR) no próprio acelerador.
    • A palavra-chave restringir proíbe o otimizador de usar referências indiretas de buffer.
    • No sexto parâmetro, um ponteiro para a paleta é passado.
    • A palavra-chave __constant otimiza os acessos ao buffer, gerando um cache com um atributo somente leitura.

    A descrição da função na listagem está próxima da implementação do processador x64. Aqui, a definição da instância atual do kernel é feita através da função get_global_id, na qual o número da dimensão (0, 1) é passado como parâmetro.

    Para uma melhor otimização, uma indicação explícita do início do ciclo foi introduzida. Na ausência de informações sobre o número de iterações no momento da compilação, o número de etapas do loop é indicado explicitamente, pois seus próprios blocos de hardware serão criados para eles. Com esse tipo de codificação, deve-se "olhar para trás" na capacidade de um chip específico instalado no acelerador, devido ao consumo de recursos do FPGA por um número maior de ciclos.

     //////////////////////////////////////////////////////////////////// // mandelbrot_kernel.cl : Hardware implementation of the mandelbrot algorithm //////////////////////////////////////////////////////////////////// // Amount of loop unrolling. #ifndef UNROLL #define UNROLL 20 #endif // Define the color black as 0 #define BLACK 0x00000000 __kernel void hw_mandelbrot_frame ( const double x0, const double y0, const double stepSize, const unsigned int maxIterations, __global unsigned int *restrict framebuffer, __constant const unsigned int *restrict colorLUT, const unsigned int windowWidth) { // Work-item position const size_t windowPosX = get_global_id(0); const size_t windowPosY = get_global_id(1); const double stepPosX = x0 + (windowPosX * stepSize); const double stepPosY = y0 - (windowPosY * stepSize); // Variables for the calculation double x = 0.0; double y = 0.0; double xSqr = 0.0; double ySqr = 0.0;</code> <code>unsigned #pragma while { int iterations = 0; // Perform up to the maximum number of iterations to solve // the current work-item's position in the image // The loop unrolling factor can be adjusted based on the amount of FPGA // resources available. unroll UNROLL xSqr + ySqr < 4.0 && iterations < maxIterations ) // Perform the current iteration xSqr = x*x; ySqr = y*y; y = 2*x*y + stepPosY; x = xSqr - ySqr + stepPosX; // Increment iteration count iterations++; } // Output black if we never finished, and a color from the look up table otherwise framebuffer[windowWidth * windowPosY + windowPosX] = (iterations == maxIterations) ? BLACK : colorLUT[iterations]; }

    O pacote de utilitários Intel FPGA SDK para OpenCL precisará ser instalado no host antes de compilar a implementação de hardware do algoritmo. Entre o software pré-instalado, você deve incluir o BSP (Board Support Package) do fabricante da placa aceleradora específica. No exemplo, o Intel Quartus Prime Pro 16.1 é instalado com suporte para OpenCL e BSP do acelerador Euler Thread (Intel Arria 10).

    A seguir, é apresentada a configuração de caminhos e variáveis ​​de ambiente. A variável ALTERAOCLSDKROOT contém o caminho para o Intel FPGA SDK, a variável AOCL_BOARD_PACKAGE_ROOT contém o caminho para o acelerador BSP.

     set ALTERAOCLSDKROOT=C:\intelFPGA_pro\16.1\hld set AOCL_BOARD_PACKAGE_ROOT=C:\intelFPGA_pro\16.1\hld\board\euler_thread set path=%path%;C:\intelFPGA_pro\16.1\hld\bin set path=%path%;C:\intelFPGA_pro\16.1\quartus\bin64 set path=%path%;C:\intelFPGA_pro\16.1\hld\board\a10_ref\windows64\bin set path=%path%;C:\intelFPGA_pro\16.1\hld\host\windows64\bin set path=%path%;C:\intelFPGA_pro\16.1\qsys\bin set path=%path%;C:\Program Files (x86)\GnuWin32\bin\

    Para compilação, o compilador aoc do SDK é usado.

     aoc mandelbrot_kernel.cl -o mandelbrot_kernel.aocx --board thread -v -v --report

    Descriptografamos: mandelbrot_kernel.cl - o arquivo com o texto de origem, mandelbrot_kernel.aocx - o arquivo de objeto de saída para programação de FPGA, thread - o nome do acelerador do pacote BSP. A opção --report exibe um relatório de uso de recursos do FPGA. A opção –v exibe informações de diagnóstico durante a compilação. O relatório de consumo de recursos para o kernel é o seguinte:

    + ------------------------------------------------- ------------------- +
    ; Resumo de uso estimado de recursos;
    + ---------------------------------------- + -------- ------------------- +
    ; Recurso + Uso;
    + ---------------------------------------- + -------- ------------------- +
    ; Utilização lógica; 49%
    ; ALUTs; 26%;
    ; Registradores lógicos dedicados; 25%;
    ; Blocos de memória; 21%
    ; Blocos DSP; 16%;
    + ---------------------------------------- + -------- -------------------;

    Para compilar o aplicativo host, o exemplo usou o pacote Microsoft Visual Studio 2010 Express com o Microsoft SDK 7.1 instalado. Nas configurações do projeto, a configuração para x64 é selecionada. Em seguida, conecte a pasta para arquivos de cabeçalho externos e especifique o caminho para bibliotecas adicionais do Intel FPGA SDK nas configurações do vinculador.
    Diretórios adicionais para incluir arquivos = $ (ALTERAOCLSDKROOT) \ host \ include;
    Diretórios adicionais da biblioteca = $ (AOCL_BOARD_PACKAGE_ROOT) \ windows64 \ lib;
     $(ALTERAOCLSDKROOT)\host\windows64\lib;

    O plano de ação geral para iniciar o kernel no acelerador será o seguinte:

    1. Obtenha uma lista de plataformas
    2. Obter uma lista de dispositivos
    3. criar contexto;
    4. carregar o kernel no dispositivo;
    5. envie buffers de entrada para o dispositivo;
    6. execute o kernel para execução;
    7. leia o buffer de saída do dispositivo;
    8. contexto livre.

    Considere alguns pontos diretamente relacionados ao lançamento do kernel. Portanto, um núcleo é projetado para processar um pixel da imagem. Portanto, você precisa executar N instâncias do kernel, em que N é o número total de pixels na imagem.

    Abaixo, observamos o caso em que existem várias placas aceleradoras no servidor e a tarefa pode ser distribuída entre elas. Em cada um dos aceleradores, você precisa carregar o kernel (arquivo mandelbrot_kernel.aocx). Suponha que o número de aceleradores seja numDevices e as linhas da imagem sejam divididas entre todos os aceleradores:

     #define MAXDEV 10 static cl_context theContext; static cl_program theProgram; static cl_kernel theKernels[MAXDEV]; //.. // Create the program object theProgram = createProgramFromBinary( theContext, "mandelbrot_kernel.aocx", theDevices, numDevices); // Create the kernels for ( unsigned i = 0; i < numDevices; ++i ) theKernels[i] = clCreateKernel( theProgram, "hw_mandelbrot_frame", &theStatus ); // Create output pixel buffers for every kernel for( unsigned i = 0; i < numDevices; ++i ) thePixelData[i] = clCreateBuffer(theContext, CL_MEM_WRITE_ONLY, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned int), NULL, &theStatus); // Preparing and writing palette buffer to every device theHardColorTable = clCreateBuffer(theContext, CL_MEM_READ_ONLY, aColorTableSize*sizeof(unsigned int), NULL, &theStatus); for( unsigned i = 0; i < numDevices; i++ ) theStatus = clEnqueueWriteBuffer(theQueues[i], theHardColorTable, CL_TRUE, 0, aColorTableSize*sizeof(unsigned int), aColorTable, 0, NULL, NULL); // Preparing kernels and run unsigned rowOffset = 0; for ( unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++] ) { // Create ND range size size_t globalSize[2] = { thePixelDataWidth, rowsPerDevice[i] }; // Set the arguments unsigned argi = 0; theStatus = clSetKernelArg (theKernels[i], argi++, sizeof(cl_double), (void*) &aStartX ); const double offsetedStartY = aStartY - rowOffset * aScale; theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_double), (void*)&offsetedStartY); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_double), (void*)&aScale); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theHardColorTableSize); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&thePixelData[i]); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_mem), (void*)&theHardColorTable); theStatus = clSetKernelArg(theKernels[i], argi++, sizeof(cl_uint), (void*)&theWidth); // Launch kernel theStatus = clEnqueueNDRangeKernel(theQueues[i], theKernels[i], 2, NULL, globalSize, NULL, 0, NULL, NULL); } rowOffset = 0; for( unsigned i = 0; i < numDevices; rowOffset += rowsPerDevice[i++] ) { // Read the output theStatus = clEnqueueReadBuffer(theQueues[i], thePixelData[i], CL_TRUE, 0, thePixelDataWidth*rowsPerDevice[i]*sizeof(unsigned int), &aFrameBuffer[rowOffset * theWidth], 0, NULL, NULL); } / / . .

    • A função createProgramFromBinary cria um objeto de programa OpenCL a partir de um arquivo de objeto.
    • Em seguida, para cada dispositivo, um kernel é criado com base no objeto do programa.
    • Os buffersPixelData são criados para receber a saída de cada núcleo.
    • Um buffer é criado para armazenar a paleta de cores e carregado em cada um dos aceleradores.
    • Em seguida, para cada dispositivo, a ligação dos parâmetros locais do aplicativo e dos parâmetros do kernel é definida usando a função clSetKernelArg.
    • Os parâmetros são determinados pelos números de série na declaração da função do kernel, começando do zero.

    O próximo ponto importante é determinar o tamanho da tarefa com base no espaço do índice de acordo com a matriz globalSize. Essa matriz pode ser unidimensional, bidimensional ou tridimensional. Para cada dimensão, uma dimensão é fornecida como um número inteiro. A dimensão do espaço determinará a ordem do índice do item de trabalho no kernel.

    No exemplo, para cada núcleo, é especificado um espaço bidimensional, em que um dos eixos são os elementos da linha de pixels, o segundo é o conjunto de linhas de imagem processadas neste dispositivo. No código do kernel, o número de pixels na linha é obtido chamando get_global_id (0), o número da linha é get_global_id (1). A variável globalSize é passada para a função clEnqueueNDRangeKernel para iniciar o número necessário de instâncias do kernel a serem executadas.

    Após a conclusão da execução dos núcleos, os buffers de pixel são lidos do dispositivo para matrizes locais. Vamos avaliar o desempenho pelo número de quadros por segundo - o resultado é visível na demonstração realizada na conferência SelectelTechDay ( consulte o início do artigo ).

    Conclusão


    A programação de aceleradores FPGA em uma linguagem de alto nível sem dúvida reduziu o limiar para os desenvolvedores acessarem essa tecnologia. Por exemplo, para aqueles que estão apenas dominando este kit de ferramentas, há até uma implementação em FPGA do famoso exemplo "Hello World" .

    Mas não é tão simples. Escrever - e especialmente - depurar um algoritmo claramente funcional de um problema real aplicado ainda requer alto profissionalismo. Outra limitação é que cada chip FPGA pode executar apenas uma tarefa computacional dentro do aplicativo. Para outra tarefa, ela deve ser reprogramada novamente.
    A propósito, o modelo de uso da plataforma permite que você tenha mais de um acelerador FPGA no host, embora essa seja uma solução bastante cara.
    O host (aplicativo host) gerencia o processo de criação do contexto (estrutura de dados para o acelerador) e a fila de comandos. I.e. Um único aplicativo host, no qual existem várias subtarefas para computação paralela no FPGA, pode carregá-las em diferentes aceleradores:
    KERNEL1 => ACELERADOR A
    KERNEL2 => ACELERADOR B

    No entanto, os esforços para dominar os aceleradores FPGA valem a pena - em muitas áreas de aplicação, essa tecnologia está se tornando indispensável: telecomunicações, biotecnologia, processamento de big data, reconhecimento de padrões, processamento de sinais e imagens, em matemática computacional e modelagem de campos físicos.

    Informações adicionais para o artigo:
    www.altera.com é o principal recurso do Intel FPGA.
    www.eulerproject.com é o site oficial do Projeto Euler.
    Altera + OpenCL: programamos em FPGA sem o conhecimento de VHDL / Verilog - um artigo sobre Habr.

Source: https://habr.com/ru/post/pt418403/

More articles:


All Articles