Equação de Navier-Stokes e simulação de fluidos na CUDA

Oi Habr. Neste artigo, trataremos da equação de Navier-Stokes para um fluido incompressível, resolveremos numericamente e faremos uma bela simulação que funciona por computação paralela na CUDA. O objetivo principal é mostrar como você pode aplicar a matemática subjacente à equação na prática ao resolver o problema de modelagem de líquidos e gases.



Advertência
O artigo contém muita matemática; portanto, aqueles que estão interessados ​​no lado técnico da questão podem ir diretamente para a seção de implementação do algoritmo. No entanto, eu ainda recomendo que você leia o artigo inteiro e tente entender o princípio da solução. Se você tiver alguma dúvida até o final da leitura, terei o maior prazer em respondê-la nos comentários do post.

Nota: se você estiver lendo o Habr em um dispositivo móvel e não vir as fórmulas, use a versão completa do site

Equação de Navier-Stokes para um fluido incompressível


 parcial bf vecu over tparcial=( bf vecu cdot nabla) bf vecu1 over rho nabla bfp+ nu nabla2 bf vecu+ bf vecF


Acho que todos pelo menos uma vez ouviram falar dessa equação, alguns, talvez até analiticamente, resolvam seus casos particulares, mas em termos gerais esse problema permanece por resolver até agora. Obviamente, não definimos o objetivo de resolver o problema do milênio neste artigo, mas ainda podemos aplicar o método iterativo a ele. Mas, para começar, vejamos a notação nesta fórmula.

Convencionalmente, a equação de Navier-Stokes pode ser dividida em cinco partes:

  •  parcial bf vecu mais parcialt - denota a taxa de variação da velocidade do fluido em um ponto (a consideraremos para cada partícula em nossa simulação).
  • ( bf vecu cdot nabla) bf vecu - o movimento do fluido no espaço.
  • 1 over rho nabla bfp A pressão exercida sobre a partícula (aqui  rho - coeficiente de densidade do fluido).
  •  nu nabla2 bf vecu - viscosidade do meio (quanto maior, mais forte o líquido resiste à força aplicada à sua parte),  nu - coeficiente de viscosidade).
  •  bf vecF - forças externas que aplicamos ao fluido (no nosso caso, a força terá um papel muito específico - refletirá as ações executadas pelo usuário.

Além disso, como consideraremos o caso de um fluido incompressível e homogêneo, temos outra equação:   nabla cdot bf vecu=0 . A energia no ambiente é constante, não vai a lugar nenhum, não vem de nenhum lugar.

Será errado privar todos os leitores que não estão familiarizados com a análise vetorial , portanto, ao mesmo tempo, e passar brevemente por todos os operadores presentes na equação (no entanto, eu recomendo fortemente lembrar qual é a derivada, o diferencial e o vetor, pois eles sustentam tudo isso. o que será discutido abaixo).

Começamos com o operador nabla, que é esse vetor (no nosso caso, será de dois componentes, pois modelaremos o fluido no espaço bidimensional):

 nabla= beginpmatrix parcial sob parcialx, parcial sobre parcialy endpmatrix


O operador nabla é um operador diferencial de vetor e pode ser aplicado tanto a uma função escalar quanto a uma função vetorial. No caso de um escalar, obtemos o gradiente da função (o vetor de suas derivadas parciais) e, no caso de um vetor, a soma das derivadas parciais ao longo dos eixos. A principal característica desse operador é que, através dele, você pode expressar as principais operações da análise vetorial - grad ( gradiente ), div ( divergência ), rot ( rotor ) e  nabla2 ( Operador Laplace ). Vale a pena notar imediatamente que a expressão ( bf vecu cdot nabla) bf vecu não equivale a ( nabla cdot bf vecu) bf vecu - o operador nabla não possui comutatividade.

Como veremos mais adiante, essas expressões são visivelmente simplificadas ao mudar para um espaço discreto no qual realizaremos todos os cálculos; portanto, não fique alarmado se no momento você não estiver muito claro sobre o que fazer com tudo isso. Depois de dividir a tarefa em várias partes, resolveremos sucessivamente cada uma delas e apresentaremos tudo isso na forma da aplicação seqüencial de várias funções em nosso ambiente.

Solução numérica da equação de Navier-Stokes


Para representar nosso fluido no programa, precisamos obter uma representação matemática do estado de cada partícula de fluido em um ponto arbitrário no tempo. O método mais conveniente para isso é criar um campo vetorial de partículas armazenando seu estado na forma de um plano de coordenadas:

imagem

Em cada célula da nossa matriz bidimensional, armazenaremos a velocidade das partículas por vez t: bf vecu=u( bf vecx,t), bf vecx= beginpmatrixx,y endpmatrix , e a distância entre as partículas é indicada por  deltax e  deltay em conformidade. No código, será suficiente alterar o valor da velocidade a cada iteração, resolvendo um conjunto de várias equações.

Agora, expressamos o gradiente, a divergência e o operador Laplace, levando em consideração nossa grade de coordenadas ( i,j - índices na matriz,  bf vecu(x), vecu(y) - retirando os componentes correspondentes do vetor):
OperadorDefinição deAnalógico Discreto
grad nabla bfp= beginpmatrix parcial bfp over parcialx, parcial bfp over parcialy endpmatrixpi+1,jpi1,j sobre2 deltax,pi,j+1pi,j1 over2 deltay
div nabla cdot bf vecu= uparcial over parcialx+ uparcial over parcialy  vecu(x)i+1,j vecu(x)i1,j sobre2 deltax+ vecu(y)i,j+1 vecu(y)i,j1 maisde2 deltay
 bf Delta bf nabla2p= parcial2p sobre parcialx2+ parcial2p sobre parcialy2pi+1,j2pi,j+pi1,j over( deltax)2+pi,j+12pi,j+pi,j1 over( deltay)2
apodrecer bf nabla times vecu= parcial v u sobre parcialy parcial vecu sobre parcialx vecu(y)i,j+1 vecu(y)i,j1 maisde2 deltay vecu(x)i+1,j vecu(x)i1,j maisde2 deltax

Podemos simplificar ainda mais as fórmulas discretas dos operadores de vetores se assumirmos que  deltax= deltay=1 . Essa suposição não afetará muito a precisão do algoritmo, no entanto, reduz o número de operações por iteração e geralmente torna as expressões mais agradáveis ​​aos olhos.

Movimento de partículas


Essas declarações funcionam apenas se pudermos encontrar as partículas mais próximas em relação à que está sendo considerada no momento. Para anular todos os custos possíveis associados à sua localização, rastrearemos não o movimento deles, mas de onde as partículas vêm no início da iteração, projetando a trajetória do movimento para trás no tempo (em outras palavras, subtraia o vetor de velocidade vezes a mudança de tempo de posição atual). Usando esta técnica para cada elemento da matriz, teremos certeza de que qualquer partícula terá "vizinhos":

imagem

Colocando isso q - um elemento da matriz que armazena o estado da partícula, obtemos a seguinte fórmula para calcular seu estado ao longo do tempo  deltat (acreditamos que todos os parâmetros necessários na forma de aceleração e pressão já foram calculados):

q(  vec bfx,t+ deltat)=q(  bf vecx bf vecu deltat,t)


Notamos imediatamente que, para pequenas quantidades  deltat e nunca podemos ir além dos limites da célula, portanto, é muito importante escolher o momento certo que o usuário dará às partículas.

Para evitar a perda de precisão no caso de uma projeção no limite da célula ou no caso de coordenadas não inteiras, realizaremos a interpolação bilinear dos estados das quatro partículas mais próximas e a consideraremos como o valor verdadeiro em um ponto. Em princípio, esse método praticamente não reduz a precisão da simulação e, ao mesmo tempo, é bastante simples de implementar, por isso vamos usá-lo.

Viscosidade


Cada líquido tem uma certa viscosidade, a capacidade de impedir a influência de forças externas em suas partes (mel e água serão um bom exemplo, em alguns casos, seus coeficientes de viscosidade diferem em uma ordem de magnitude). A viscosidade afeta diretamente a aceleração adquirida pelo líquido e pode ser expressa pela fórmula abaixo, se, por questões de brevidade, omitirmos outros termos por um tempo:

 parcial vec bfu sobre parcialt= nu nabla2 bf vecu

. Nesse caso, a equação iterativa para velocidade assume a seguinte forma:

u( bf vecx,t+ deltat)=u( bf vecx,t)+ nu deltat nabla2 bf vecu


Transformaremos levemente essa igualdade, trazendo-a para a forma  bfA vecx= vecb (forma padrão de um sistema de equações lineares):

( bfI nu deltat nabla2)u( bf vecx,t+ deltat)=u( bf vecx,t)


onde  bfI É a matriz de identidade. Precisamos dessas transformações para posteriormente aplicar o método de Jacobi para resolver vários sistemas de equações semelhantes. Também discutiremos isso mais tarde.

Forças externas


O passo mais simples do algoritmo é a aplicação de forças externas ao meio. Para o usuário, isso será refletido na forma de cliques na tela com o mouse ou seu movimento. A força externa pode ser descrita pela fórmula a seguir, que aplicamos a cada elemento da matriz (  vec bfG - vetor de momento xp,yp - posição do mouse x,y - coordenadas da célula atual, r - raio de ação, parâmetro de escala):

 vec bfF= vec bfG deltat bfexp left((xxp)2+(yyp)2 sobrer right)


Um vetor de impulso pode ser facilmente calculado como a diferença entre a posição anterior do mouse e a atual (se houver), e aqui você ainda pode ser criativo. É nesta parte do algoritmo que podemos introduzir a adição de cores a um líquido, sua iluminação etc. As forças externas também podem incluir gravidade e temperatura e, embora não seja difícil implementar tais parâmetros, não os consideraremos neste artigo.

Pressão


A pressão na equação de Navier-Stokes é a força que impede que as partículas preencham todo o espaço disponível após aplicar qualquer força externa nelas. Imediatamente, seu cálculo é muito difícil, mas nosso problema pode ser bastante simplificado aplicando o teorema da decomposição de Helmholtz .

Ligar  bf vecW campo vetorial obtido após o cálculo do deslocamento, forças externas e viscosidade. Apresentará uma divergência diferente de zero, o que contradiz a condição de incompressibilidade do líquido (  nabla cdot bf vecu=0 ) e, para corrigir isso, é necessário calcular a pressão. De acordo com o teorema da decomposição de Helmholtz,  bf vecW pode ser representado como a soma de dois campos:

 bf vecW= vecu+ nap


onde  bfu - este é o campo vetorial que procuramos com divergência zero. Nenhuma prova dessa igualdade será fornecida neste artigo, mas no final você pode encontrar um link com uma explicação detalhada. Podemos aplicar o operador nabla nos dois lados da expressão para obter a seguinte fórmula para calcular o campo de pressão escalar:

 nabla cdot bf vecW= nabla cdot( vecu+ nablap)= nabla cdot vecu+ nabla2p= nabla2p


A expressão escrita acima é a equação de Poisson para pressão. Também podemos resolvê-lo pelo método de Jacobi acima mencionado e, assim, encontrar a última variável desconhecida na equação de Navier-Stokes. Em princípio, os sistemas de equações lineares podem ser resolvidos de várias maneiras diferentes e sofisticadas, mas ainda vamos nos concentrar nas mais simples delas, para não sobrecarregar ainda mais este artigo.

Limite e condições iniciais


Qualquer equação diferencial modelada em um domínio finito requer condições iniciais ou de contorno especificadas corretamente, caso contrário, é muito provável que obtenhamos um resultado fisicamente incorreto. As condições de contorno são estabelecidas para controlar o comportamento do fluido próximo às bordas da grade de coordenadas, e as condições iniciais especificam os parâmetros que as partículas têm no momento em que o programa é iniciado.

As condições iniciais serão muito simples - inicialmente o fluido é estacionário (a velocidade das partículas é zero) e a pressão também é zero. As condições de contorno serão definidas para velocidade e pressão pelas fórmulas fornecidas:

 bf vecu0,j+ bf vecu1,j maisde2 deltay=0, bf vecui,0+ bf vecui,1 maisde2 deltax=0

 bfp0,j bfp1,j over deltax=0, bfpi,0 bfpi,1 over deltay=0


Assim, a velocidade das partículas nas bordas será oposta à velocidade nas bordas (assim elas se repelirão a partir da borda), e a pressão é igual ao valor imediatamente próximo ao limite. Essas operações devem ser aplicadas a todos os elementos delimitadores da matriz (por exemplo, há um tamanho de grade N vezesM , aplicamos o algoritmo às células marcadas em azul na figura):

imagem

Corante


Com o que temos agora, você já pode ter muitas coisas interessantes. Por exemplo, para perceber a propagação do corante em um líquido. Para fazer isso, precisamos apenas manter outro campo escalar, que seria responsável pela quantidade de tinta em cada ponto da simulação. A fórmula para atualizar o corante é muito semelhante à velocidade e é expressa como:

 parciald sobre parcialt=( vec bfu cdot nabla)d+ gama nabla2d+S


Na fórmula S responsável por reabastecer a área com um corante (possivelmente dependendo de onde o usuário clica), d diretamente é a quantidade de corante no ponto e  gama - coeficiente de difusão. Não é difícil resolvê-lo, pois todo o trabalho básico sobre derivação de fórmulas já foi realizado e basta fazer algumas substituições. A tinta pode ser implementada no código como uma cor no formato RGB e, nesse caso, a tarefa é reduzida a operações com vários valores reais.

Vorticidade


A equação de vorticidade não é uma parte direta da equação de Navier-Stokes, mas é um parâmetro importante para uma simulação plausível do movimento de um corante em um líquido. Devido ao fato de estarmos produzindo um algoritmo em um campo discreto, bem como devido a perdas na precisão dos valores de ponto flutuante, esse efeito é perdido e, portanto, precisamos restaurá-lo aplicando força adicional a cada ponto no espaço. O vetor dessa força é designado como  bf vecT e é determinado pelas seguintes fórmulas:

 bf omega= nabla times vecu

 vec eta= nabla| omega|

 bf vec psi= vec eta sobre| vec eta|

 bf vecT= epsilon( vec psi times omega) deltax


 omega há o resultado da aplicação do rotor ao vetor de velocidade (sua definição é dada no início do artigo),  vec eta - gradiente do campo escalar de valores absolutos  omega .  vec psi representa um vetor normalizado  vec eta e  epsilon É uma constante que controla o tamanho dos vórtices em nosso fluido.

Método de Jacobi para resolver sistemas de equações lineares


Ao analisar as equações de Navier-Stokes, criamos dois sistemas de equações - um para viscosidade e outro para pressão. Eles podem ser resolvidos por um algoritmo iterativo, que pode ser descrito pela seguinte fórmula iterativa:

x(k+1)i,j=x(k)i1,j+x(k)i+1,j+x(k)i,j1+x(k)i,j+1+ alphabi,j over beta


Para nós x - elementos de matriz que representam um campo escalar ou vetorial. k - o número da iteração, podemos ajustá-lo para aumentar a precisão do cálculo ou vice-versa para reduzi-lo e aumentar a produtividade.

Para calcular a viscosidade, substituímos: x=b= bf vecu ,  alpha=1 over nu deltat ,  beta=4+ alpha , aqui está o parâmetro  beta - a soma dos pesos. Portanto, precisamos armazenar pelo menos dois campos de velocidade vetorial para ler independentemente os valores de um campo e gravá-los em outro. Em média, para calcular o campo de velocidade pelo método Jacobi, é necessário realizar de 20 a 50 iterações, o que é bastante se realizarmos cálculos na CPU.

Para a equação da pressão, fazemos a seguinte substituição: x=p , b= nabla bf cdot vecW ,  a l p h a = - 1 ,  b e t a = 4 . Como resultado, obtemos o valor p i , j d e l t a t  no ponto. Porém, como é usado apenas para calcular o gradiente subtraído do campo de velocidade, transformações adicionais podem ser omitidas. Para o campo de pressão, é melhor executar iterações de 40 a 80, porque com números menores a discrepância se torna perceptível.

Implementação de algoritmo


Vamos implementar o algoritmo em C ++, também precisamos do Cuda Toolkit (você pode ler como instalá-lo no site da Nvidia), bem como do SFML . Precisamos que o CUDA paralelize o algoritmo, e o SFML será usado apenas para criar uma janela e exibir uma imagem na tela (em princípio, isso pode ser escrito em OpenGL, mas a diferença de desempenho não será significativa, mas o código aumentará em mais de 200 linhas).

Cuda toolkit


Primeiro, falaremos um pouco sobre como usar o Cuda Toolkit para paralelizar tarefas. Um guia mais detalhado é fornecido pela própria Nvidia; portanto, aqui nos restringimos apenas ao mais necessário. Supõe-se também que você conseguiu instalar o compilador e conseguiu criar um projeto de teste sem erros.

Para criar uma função que é executada na GPU, primeiro você precisa declarar quantos kernels queremos usar e quantos blocos de kernels precisamos alocar. Para isso, o Cuda Toolkit nos fornece uma estrutura especial - dim3 , por padrão, configurando todos os seus valores de x, y, z como 1. Ao especificá-lo como argumento ao chamar a função, podemos controlar o número de kernels alocados. Como estamos trabalhando com uma matriz bidimensional, precisamos definir apenas dois campos no construtor: xey :

dim3 threadsPerBlock(x_threads, y_threads); dim3 numBlocks(size_x / x_threads, y_size / y_threads); 

onde size_x e size_y são o tamanho da matriz que está sendo processada. A assinatura e a chamada de função são as seguintes (colchetes de ângulo triplo são processados ​​pelo compilador Cuda):

 void __global__ deviceFunction(); // declare deviceFunction<<<numBlocks, threadsPerBlock>>>(); // call from host 

Na própria função, você pode restaurar os índices de uma matriz bidimensional através do número do bloco e do número do kernel nesse bloco, usando a seguinte fórmula:

 int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; 

Deve-se observar que a função executada na placa de vídeo deve ser marcada com a tag __global__ e também retornar nula , portanto, na maioria das vezes, os resultados do cálculo são gravados na matriz passada como argumento e pré-alocada na memória da placa de vídeo.

As funções CudaMalloc e CudaFree são responsáveis ​​por liberar e alocar memória na placa de vídeo. Podemos operar ponteiros para a área de memória que eles retornam, mas não podemos acessar os dados do código principal. A maneira mais fácil de retornar os resultados do cálculo é usar cudaMemcpy , semelhante ao memcpy padrão, mas capaz de copiar dados de uma placa de vídeo para a memória principal e vice-versa.

SFML e renderização de janela


Armado com todo esse conhecimento, podemos finalmente começar a escrever código diretamente. Para começar, vamos criar o arquivo main.cpp e colocar todo o código auxiliar para a janela renderizada lá:

main.cpp
 #include <SFML/Graphics.hpp> #include <chrono> #include <cstdlib> #include <cmath> //SFML REQUIRED TO LAUNCH THIS CODE #define SCALE 2 #define WINDOW_WIDTH 1280 #define WINDOW_HEIGHT 720 #define FIELD_WIDTH WINDOW_WIDTH / SCALE #define FIELD_HEIGHT WINDOW_HEIGHT / SCALE static struct Config { float velocityDiffusion; float pressure; float vorticity; float colorDiffusion; float densityDiffusion; float forceScale; float bloomIntense; int radius; bool bloomEnabled; } config; void setConfig(float vDiffusion = 0.8f, float pressure = 1.5f, float vorticity = 50.0f, float cDiffuion = 0.8f, float dDiffuion = 1.2f, float force = 1000.0f, float bloomIntense = 25000.0f, int radius = 100, bool bloom = true); void computeField(uint8_t* result, float dt, int x1pos = -1, int y1pos = -1, int x2pos = -1, int y2pos = -1, bool isPressed = false); void cudaInit(size_t xSize, size_t ySize); void cudaExit(); int main() { cudaInit(FIELD_WIDTH, FIELD_HEIGHT); srand(time(NULL)); sf::RenderWindow window(sf::VideoMode(WINDOW_WIDTH, WINDOW_HEIGHT), ""); auto start = std::chrono::system_clock::now(); auto end = std::chrono::system_clock::now(); sf::Texture texture; sf::Sprite sprite; std::vector<sf::Uint8> pixelBuffer(FIELD_WIDTH * FIELD_HEIGHT * 4); texture.create(FIELD_WIDTH, FIELD_HEIGHT); sf::Vector2i mpos1 = { -1, -1 }, mpos2 = { -1, -1 }; bool isPressed = false; bool isPaused = false; while (window.isOpen()) { end = std::chrono::system_clock::now(); std::chrono::duration<float> diff = end - start; window.setTitle("Fluid simulator " + std::to_string(int(1.0f / diff.count())) + " fps"); start = end; window.clear(sf::Color::White); sf::Event event; while (window.pollEvent(event)) { if (event.type == sf::Event::Closed) window.close(); if (event.type == sf::Event::MouseButtonPressed) { if (event.mouseButton.button == sf::Mouse::Button::Left) { mpos1 = { event.mouseButton.x, event.mouseButton.y }; mpos1 /= SCALE; isPressed = true; } else { isPaused = !isPaused; } } if (event.type == sf::Event::MouseButtonReleased) { isPressed = false; } if (event.type == sf::Event::MouseMoved) { std::swap(mpos1, mpos2); mpos2 = { event.mouseMove.x, event.mouseMove.y }; mpos2 /= SCALE; } } float dt = 0.02f; if (!isPaused) computeField(pixelBuffer.data(), dt, mpos1.x, mpos1.y, mpos2.x, mpos2.y, isPressed); texture.update(pixelBuffer.data()); sprite.setTexture(texture); sprite.setScale({ SCALE, SCALE }); window.draw(sprite); window.display(); } cudaExit(); return 0; } 


linha no início da função principal

 std::vector<sf::Uint8> pixelBuffer(FIELD_WIDTH * FIELD_HEIGHT * 4); 

cria uma imagem RGBA na forma de uma matriz unidimensional com um comprimento constante. Passaremos junto com outros parâmetros (posição do mouse, diferença entre os quadros) para a função computeField . As últimas, assim como várias outras funções, são declaradas no kernel.cu e chamam o código executado na GPU. Você pode encontrar documentação sobre qualquer uma das funções no site da SFML, nada super interessante está acontecendo no código do arquivo, por isso não vamos parar por muito tempo.

Computação GPU


Para começar a escrever código no gpu, primeiro crie um arquivo kernel.cu e defina várias classes auxiliares: Color3f, Vec2, Config, SystemConfig :

kernel.cu (estruturas de dados)
 struct Vec2 { float x = 0.0, y = 0.0; __device__ Vec2 operator-(Vec2 other) { Vec2 res; res.x = this->x - other.x; res.y = this->y - other.y; return res; } __device__ Vec2 operator+(Vec2 other) { Vec2 res; res.x = this->x + other.x; res.y = this->y + other.y; return res; } __device__ Vec2 operator*(float d) { Vec2 res; res.x = this->x * d; res.y = this->y * d; return res; } }; struct Color3f { float R = 0.0f; float G = 0.0f; float B = 0.0f; __host__ __device__ Color3f operator+ (Color3f other) { Color3f res; res.R = this->R + other.R; res.G = this->G + other.G; res.B = this->B + other.B; return res; } __host__ __device__ Color3f operator* (float d) { Color3f res; res.R = this->R * d; res.G = this->G * d; res.B = this->B * d; return res; } }; struct Particle { Vec2 u; // velocity Color3f color; }; static struct Config { float velocityDiffusion; float pressure; float vorticity; float colorDiffusion; float densityDiffusion; float forceScale; float bloomIntense; int radius; bool bloomEnabled; } config; static struct SystemConfig { int velocityIterations = 20; int pressureIterations = 40; int xThreads = 64; int yThreads = 1; } sConfig; void setConfig( float vDiffusion = 0.8f, float pressure = 1.5f, float vorticity = 50.0f, float cDiffuion = 0.8f, float dDiffuion = 1.2f, float force = 5000.0f, float bloomIntense = 25000.0f, int radius = 500, bool bloom = true ) { config.velocityDiffusion = vDiffusion; config.pressure = pressure; config.vorticity = vorticity; config.colorDiffusion = cDiffuion; config.densityDiffusion = dDiffuion; config.forceScale = force; config.bloomIntense = bloomIntense; config.radius = radius; config.bloomEnabled = bloom; } static const int colorArraySize = 7; Color3f colorArray[colorArraySize]; static Particle* newField; static Particle* oldField; static uint8_t* colorField; static size_t xSize, ySize; static float* pressureOld; static float* pressureNew; static float* vorticityField; static Color3f currentColor; static float elapsedTime = 0.0f; static float timeSincePress = 0.0f; static float bloomIntense; int lastXpos = -1, lastYpos = -1; 

O atributo __host__ na frente do nome do método significa que o código pode ser executado na CPU; __dispositivo__ , pelo contrário, obriga o compilador a coletar código na GPU. O código declara primitivas para trabalhar com vetores de dois componentes, cor, configurações com parâmetros que podem ser alterados em tempo de execução, além de vários ponteiros estáticos para matrizes, que usaremos como buffers para cálculos.

cudaInit e cudaExit também são definidos de maneira bastante trivial:

kernel.cu (init)
 void cudaInit(size_t x, size_t y) { setConfig(); colorArray[0] = { 1.0f, 0.0f, 0.0f }; colorArray[1] = { 0.0f, 1.0f, 0.0f }; colorArray[2] = { 1.0f, 0.0f, 1.0f }; colorArray[3] = { 1.0f, 1.0f, 0.0f }; colorArray[4] = { 0.0f, 1.0f, 1.0f }; colorArray[5] = { 1.0f, 0.0f, 1.0f }; colorArray[6] = { 1.0f, 0.5f, 0.3f }; int idx = rand() % colorArraySize; currentColor = colorArray[idx]; xSize = x, ySize = y; cudaSetDevice(0); cudaMalloc(&colorField, xSize * ySize * 4 * sizeof(uint8_t)); cudaMalloc(&oldField, xSize * ySize * sizeof(Particle)); cudaMalloc(&newField, xSize * ySize * sizeof(Particle)); cudaMalloc(&pressureOld, xSize * ySize * sizeof(float)); cudaMalloc(&pressureNew, xSize * ySize * sizeof(float)); cudaMalloc(&vorticityField, xSize * ySize * sizeof(float)); } void cudaExit() { cudaFree(colorField); cudaFree(oldField); cudaFree(newField); cudaFree(pressureOld); cudaFree(pressureNew); cudaFree(vorticityField); } 

Na função de inicialização, alocamos memória para matrizes bidimensionais, especificamos uma matriz de cores que usaremos para pintar o líquido e também definimos os valores padrão na configuração. Em cudaExit, apenas liberamos todos os buffers. Por mais paradoxal que pareça, para armazenar matrizes bidimensionais, é mais vantajoso usar matrizes unidimensionais, cujo acesso será realizado com a seguinte expressão:

 array[y * size_x + x]; // equals to array[y][x] 

Iniciamos a implementação do algoritmo direto com a função de deslocamento de partículas. Os campos oldField e newField (o campo de onde os dados vêm e para onde são gravados), o tamanho da matriz, bem como o delta de tempo e o coeficiente de densidade (usado para acelerar a dissolução do corante no líquido e tornar o meio não muito sensível a ações do usuário). A função de interpolação bilinear é implementada de maneira clássica, calculando valores intermediários:

kernel.cu (advect)
 // interpolates quantity of grid cells __device__ Particle interpolate(Vec2 v, Particle* field, size_t xSize, size_t ySize) { float x1 = (int)vx; float y1 = (int)vy; float x2 = (int)vx + 1; float y2 = (int)vy + 1; Particle q1, q2, q3, q4; #define CLAMP(val, minv, maxv) min(maxv, max(minv, val)) #define SET(Q, x, y) Q = field[int(CLAMP(y, 0.0f, ySize - 1.0f)) * xSize + int(CLAMP(x, 0.0f, xSize - 1.0f))] SET(q1, x1, y1); SET(q2, x1, y2); SET(q3, x2, y1); SET(q4, x2, y2); #undef SET #undef CLAMP float t1 = (x2 - vx) / (x2 - x1); float t2 = (vx - x1) / (x2 - x1); Vec2 f1 = q1.u * t1 + q3.u * t2; Vec2 f2 = q2.u * t1 + q4.u * t2; Color3f C1 = q2.color * t1 + q4.color * t2; Color3f C2 = q2.color * t1 + q4.color * t2; float t3 = (y2 - vy) / (y2 - y1); float t4 = (vy - y1) / (y2 - y1); Particle res; res.u = f1 * t3 + f2 * t4; res.color = C1 * t3 + C2 * t4; return res; } // adds quantity to particles using bilinear interpolation __global__ void advect(Particle* newField, Particle* oldField, size_t xSize, size_t ySize, float dDiffusion, float dt) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; float decay = 1.0f / (1.0f + dDiffusion * dt); Vec2 pos = { x * 1.0f, y * 1.0f }; Particle& Pold = oldField[y * xSize + x]; // find new particle tracing where it came from Particle p = interpolate(pos - Pold.u * dt, oldField, xSize, ySize); pu = pu * decay; p.color = p.color * decay; newField[y * xSize + x] = p; } 

Decidiu-se dividir a função de difusão da viscosidade em várias partes: computeDiffusion é chamada a partir do código principal, que chama difusa e computeColor um número predeterminado de vezes e depois troca a matriz de onde pegamos os dados e a onde os escrevemos. Essa é a maneira mais fácil de implementar o processamento de dados paralelos, mas estamos gastando o dobro de memória.

Ambas as funções causam variações do método Jacobi. O corpo de jacobiColor e jacobiVelocity verifica imediatamente se os elementos atuais não estão na borda - nesse caso, devemos defini-los de acordo com as fórmulas descritas na seção Limite e condições iniciais .

kernel.cu (difuso)
 // performs iteration of jacobi method on color grid field __device__ Color3f jacobiColor(Particle* colorField, size_t xSize, size_t ySize, Vec2 pos, Color3f B, float alpha, float beta) { Color3f xU, xD, xL, xR, res; int x = pos.x; int y = pos.y; #define SET(P, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) P = colorField[int(y) * xSize + int(x)] SET(xU, x, y - 1).color; SET(xD, x, y + 1).color; SET(xL, x - 1, y).color; SET(xR, x + 1, y).color; #undef SET res = (xU + xD + xL + xR + B * alpha) * (1.0f / beta); return res; } // calculates color field diffusion __global__ void computeColor(Particle* newField, Particle* oldField, size_t xSize, size_t ySize, float cDiffusion, float dt) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; Vec2 pos = { x * 1.0f, y * 1.0f }; Color3f c = oldField[y * xSize + x].color; float alpha = cDiffusion * cDiffusion / dt; float beta = 4.0f + alpha; // perfom one iteration of jacobi method (diffuse method should be called 20-50 times per cell) newField[y * xSize + x].color = jacobiColor(oldField, xSize, ySize, pos, c, alpha, beta); } // performs iteration of jacobi method on velocity grid field __device__ Vec2 jacobiVelocity(Particle* field, size_t xSize, size_t ySize, Vec2 v, Vec2 B, float alpha, float beta) { Vec2 vU = B * -1.0f, vD = B * -1.0f, vR = B * -1.0f, vL = B * -1.0f; #define SET(U, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) U = field[int(y) * xSize + int(x)].u SET(vU, vx, vy - 1); SET(vD, vx, vy + 1); SET(vL, vx - 1, vy); SET(vR, vx + 1, vy); #undef SET v = (vU + vD + vL + vR + B * alpha) * (1.0f / beta); return v; } // calculates nonzero divergency velocity field u __global__ void diffuse(Particle* newField, Particle* oldField, size_t xSize, size_t ySize, float vDiffusion, float dt) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; Vec2 pos = { x * 1.0f, y * 1.0f }; Vec2 u = oldField[y * xSize + x].u; // perfoms one iteration of jacobi method (diffuse method should be called 20-50 times per cell) float alpha = vDiffusion * vDiffusion / dt; float beta = 4.0f + alpha; newField[y * xSize + x].u = jacobiVelocity(oldField, xSize, ySize, pos, u, alpha, beta); } // performs several iterations over velocity and color fields void computeDiffusion(dim3 numBlocks, dim3 threadsPerBlock, float dt) { // diffuse velocity and color for (int i = 0; i < sConfig.velocityIterations; i++) { diffuse<<<numBlocks, threadsPerBlock>>>(newField, oldField, xSize, ySize, config.velocityDiffusion, dt); computeColor<<<numBlocks, threadsPerBlock>>>(newField, oldField, xSize, ySize, config.colorDiffusion, dt); std::swap(newField, oldField); } } 

O uso da força externa é implementado através de uma única função - applyForce , que leva como argumento a posição do mouse, a cor do corante e também o raio de ação. Com sua ajuda, podemos dar velocidade às partículas e pintá-las. O expoente fraterno permite que você deixe a área não muito nítida e, ao mesmo tempo, bastante clara no raio especificado.

kernel.cu (força)
 // applies force and add color dye to the particle field __global__ void applyForce(Particle* field, size_t xSize, size_t ySize, Color3f color, Vec2 F, Vec2 pos, int r, float dt) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; float e = expf((-(powf(x - pos.x, 2) + powf(y - pos.y, 2))) / r); Vec2 uF = F * dt * e; Particle& p = field[y * xSize + x]; pu = pu + uF; color = color * e + p.color; p.color.R = min(1.0f, color.R); p.color.G = min(1.0f, color.G); p.color.B = min(1.0f, color.B); } 

O cálculo da vorticidade já é um processo mais complicado, portanto, implementado em computateVorticity e applyVorticity , também observamos que para eles é necessário definir dois operadores vetoriais como curl (rotor) e absGradient (gradiente de valores absolutos de campo). Para especificar efeitos de vórtice adicionais, multiplicamos y componente do vetor gradiente em - 1 e normalize-o dividindo pelo comprimento (sem esquecer de verificar se o vetor é diferente de zero):

kernel.cu (vorticidade)
 // computes curl of velocity field __device__ float curl(Particle* field, size_t xSize, size_t ySize, int x, int y) { Vec2 C = field[int(y) * xSize + int(x)].u; #define SET(P, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) P = field[int(y) * xSize + int(x)] float x1 = -Cx, x2 = -Cx, y1 = -Cy, y2 = -Cy; SET(x1, x + 1, y).ux; SET(x2, x - 1, y).ux; SET(y1, x, y + 1).uy; SET(y2, x, y - 1).uy; #undef SET float res = ((y1 - y2) - (x1 - x2)) * 0.5f; return res; } // computes absolute value gradient of vorticity field __device__ Vec2 absGradient(float* field, size_t xSize, size_t ySize, int x, int y) { float C = field[int(y) * xSize + int(x)]; #define SET(P, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) P = field[int(y) * xSize + int(x)] float x1 = C, x2 = C, y1 = C, y2 = C; SET(x1, x + 1, y); SET(x2, x - 1, y); SET(y1, x, y + 1); SET(y2, x, y - 1); #undef SET Vec2 res = { (abs(x1) - abs(x2)) * 0.5f, (abs(y1) - abs(y2)) * 0.5f }; return res; } // computes vorticity field which should be passed to applyVorticity function __global__ void computeVorticity(float* vField, Particle* field, size_t xSize, size_t ySize) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; vField[y * xSize + x] = curl(field, xSize, ySize, x, y); } // applies vorticity to velocity field __global__ void applyVorticity(Particle* newField, Particle* oldField, float* vField, size_t xSize, size_t ySize, float vorticity, float dt) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; Particle& pOld = oldField[y * xSize + x]; Particle& pNew = newField[y * xSize + x]; Vec2 v = absGradient(vField, xSize, ySize, x, y); vy *= -1.0f; float length = sqrtf(vx * vx + vy * vy) + 1e-5f; Vec2 vNorm = v * (1.0f / length); Vec2 vF = vNorm * vField[y * xSize + x] * vorticity; pNew = pOld; pNew.u = pNew.u + vF * dt; } 

O próximo passo no algoritmo será o cálculo do campo de pressão escalar e sua projeção no campo de velocidade. Para fazer isso, precisamos implementar 4 funções: divergência , que considerará a divergência de velocidade, jacobiPressure , que implementa o método Jacobi para pressão, e compututePressure com computePressureImpl , executando cálculos de campo iterativos:

kernel.cu (pressão)
 // performs iteration of jacobi method on pressure grid field __device__ float jacobiPressure(float* pressureField, size_t xSize, size_t ySize, int x, int y, float B, float alpha, float beta) { float C = pressureField[int(y) * xSize + int(x)]; float xU = C, xD = C, xL = C, xR = C; #define SET(P, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) P = pressureField[int(y) * xSize + int(x)] SET(xU, x, y - 1); SET(xD, x, y + 1); SET(xL, x - 1, y); SET(xR, x + 1, y); #undef SET float pressure = (xU + xD + xL + xR + alpha * B) * (1.0f / beta); return pressure; } // computes divergency of velocity field __device__ float divergency(Particle* field, size_t xSize, size_t ySize, int x, int y) { Particle& C = field[int(y) * xSize + int(x)]; float x1 = -1 * Cux, x2 = -1 * Cux, y1 = -1 * Cuy, y2 = -1 * Cuy; #define SET(P, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) P = field[int(y) * xSize + int(x)] SET(x1, x + 1, y).ux; SET(x2, x - 1, y).ux; SET(y1, x, y + 1).uy; SET(y2, x, y - 1).uy; #undef SET return (x1 - x2 + y1 - y2) * 0.5f; } // performs iteration of jacobi method on pressure field __global__ void computePressureImpl(Particle* field, size_t xSize, size_t ySize, float* pNew, float* pOld, float pressure, float dt) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; float div = divergency(field, xSize, ySize, x, y); float alpha = -1.0f * pressure * pressure; float beta = 4.0; pNew[y * xSize + x] = jacobiPressure(pOld, xSize, ySize, x, y, div, alpha, beta); } // performs several iterations over pressure field void computePressure(dim3 numBlocks, dim3 threadsPerBlock, float dt) { for (int i = 0; i < sConfig.pressureIterations; i++) { computePressureImpl<<<numBlocks, threadsPerBlock>>>(oldField, xSize, ySize, pressureNew, pressureOld, config.pressure, dt); std::swap(pressureOld, pressureNew); } } 

A projeção se encaixa em duas pequenas funções - projeto e o gradiente que exige pressão. Pode-se dizer o último estágio do nosso algoritmo de simulação:

kernel.cu (projeto)
 // computes gradient of pressure field __device__ Vec2 gradient(float* field, size_t xSize, size_t ySize, int x, int y) { float C = field[y * xSize + x]; #define SET(P, x, y) if (x < xSize && x >= 0 && y < ySize && y >= 0) P = field[int(y) * xSize + int(x)] float x1 = C, x2 = C, y1 = C, y2 = C; SET(x1, x + 1, y); SET(x2, x - 1, y); SET(y1, x, y + 1); SET(y2, x, y - 1); #undef SET Vec2 res = { (x1 - x2) * 0.5f, (y1 - y2) * 0.5f }; return res; } // projects pressure field on velocity field __global__ void project(Particle* newField, size_t xSize, size_t ySize, float* pField) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; Vec2& u = newField[y * xSize + x].u; u = u - gradient(pField, xSize, ySize, x, y); } 

Após a projeção, podemos proceder com segurança para renderizar a imagem no buffer e em vários pós-efeitos. A função de pintura copia cores do campo de partículas para a matriz RGBA. A função applyBloom também foi implementada, que destaca o líquido quando o cursor é colocado nele e o botão do mouse é pressionado. Por experiência, essa técnica torna a imagem mais agradável e interessante para os olhos do usuário, mas não é necessária.

No pós-processamento, você também pode destacar locais onde o fluido tem a velocidade mais alta, mudar de cor dependendo do vetor de movimento, adicionar vários efeitos etc., mas, no nosso caso, nos limitaremos a um mínimo, porque mesmo com isso as imagens são muito fascinantes (especialmente em dinâmica) :

kernel.cu (pintura)
 // adds flashlight effect near the mouse position __global__ void applyBloom(uint8_t* colorField, size_t xSize, size_t ySize, int xpos, int ypos, float bloomIntense) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; int pos = 4 * (y * xSize + x); float e = expf(-(powf(x - xpos, 2) + powf(y - ypos, 2)) * (1.0f / (bloomIntense + 1e-5f))); uint8_t R = colorField[pos + 0]; uint8_t G = colorField[pos + 1]; uint8_t B = colorField[pos + 2]; uint8_t maxval = max(R, max(G, B)); colorField[pos + 0] = min(255.0f, R + maxval * e); colorField[pos + 1] = min(255.0f, G + maxval * e); colorField[pos + 2] = min(255.0f, B + maxval * e); } // fills output image with corresponding color __global__ void paint(uint8_t* colorField, Particle* field, size_t xSize, size_t ySize) { int x = blockIdx.x * blockDim.x + threadIdx.x; int y = blockIdx.y * blockDim.y + threadIdx.y; float R = field[y * xSize + x].color.R; float G = field[y * xSize + x].color.G; float B = field[y * xSize + x].color.B; colorField[4 * (y * xSize + x) + 0] = min(255.0f, 255.0f * R); colorField[4 * (y * xSize + x) + 1] = min(255.0f, 255.0f * G); colorField[4 * (y * xSize + x) + 2] = min(255.0f, 255.0f * B); colorField[4 * (y * xSize + x) + 3] = 255; } 

E no final, ainda temos uma função principal que chamamos de main.cpp - computeField . Ele liga todas as partes do algoritmo, chamando o código na placa de vídeo, e também copia os dados de gpu para cpu. Ele também contém o cálculo do vetor de momento e a escolha da cor do corante, que passamos a aplicar a força :

kernel.cu (função principal)
 // main function, calls vorticity -> diffusion -> force -> pressure -> project -> advect -> paint -> bloom void computeField(uint8_t* result, float dt, int x1pos, int y1pos, int x2pos, int y2pos, bool isPressed) { dim3 threadsPerBlock(sConfig.xThreads, sConfig.yThreads); dim3 numBlocks(xSize / threadsPerBlock.x, ySize / threadsPerBlock.y); // curls and vortisity computeVorticity<<<numBlocks, threadsPerBlock>>>(vorticityField, oldField, xSize, ySize); applyVorticity<<<numBlocks, threadsPerBlock>>>(newField, oldField, vorticityField, xSize, ySize, config.vorticity, dt); std::swap(oldField, newField); // diffuse velocity and color computeDiffusion(numBlocks, threadsPerBlock, dt); // apply force if (isPressed) { timeSincePress = 0.0f; elapsedTime += dt; // apply gradient to color int roundT = int(elapsedTime) % colorArraySize; int ceilT = int((elapsedTime) + 1) % colorArraySize; float w = elapsedTime - int(elapsedTime); currentColor = colorArray[roundT] * (1 - w) + colorArray[ceilT] * w; Vec2 F; float scale = config.forceScale; Fx = (x2pos - x1pos) * scale; Fy = (y2pos - y1pos) * scale; Vec2 pos = { x2pos * 1.0f, y2pos * 1.0f }; applyForce<<<numBlocks, threadsPerBlock>>>(oldField, xSize, ySize, currentColor, F, pos, config.radius, dt); } else { timeSincePress += dt; } // compute pressure computePressure(numBlocks, threadsPerBlock, dt); // project project<<<numBlocks, threadsPerBlock>>>(oldField, xSize, ySize, pressureOld); cudaMemset(pressureOld, 0, xSize * ySize * sizeof(float)); // advect advect<<<numBlocks, threadsPerBlock>>>(newField, oldField, xSize, ySize, config.densityDiffusion, dt); std::swap(newField, oldField); // paint image paint<<<numBlocks, threadsPerBlock>>>(colorField, oldField, xSize, ySize); // apply bloom in mouse pos if (config.bloomEnabled && timeSincePress < 5.0f) { applyBloom<<<numBlocks, threadsPerBlock>>>(colorField, xSize, ySize, x2pos, y2pos, config.bloomIntense / timeSincePress); } // copy image to cpu size_t size = xSize * ySize * 4 * sizeof(uint8_t); cudaMemcpy(result, colorField, size, cudaMemcpyDeviceToHost); cudaError_t error = cudaGetLastError(); if (error != cudaSuccess) { std::cout << cudaGetErrorName(error) << std::endl; } } 

Conclusão


Neste artigo, analisamos um algoritmo numérico para resolver a equação de Navier-Stokes e escrevemos um pequeno programa de simulação para um fluido incompressível. Talvez não tenhamos entendido todos os meandros, mas espero que o material tenha sido interessante e útil para você, e pelo menos tenha servido como uma boa introdução ao campo da modelagem de fluidos.

Como autor deste artigo, sinceramente aprecio quaisquer comentários e adições e tentarei responder a todas as suas perguntas neste post.

Material adicional


Você pode encontrar todo o código-fonte neste artigo no meu repositório do Github . Todas as sugestões de melhoria são bem-vindas.

O material original que serviu de base para este artigo, você pode ler no site oficial da Nvidia. Também apresenta exemplos da implementação de partes do algoritmo na linguagem dos shaders:
developer.download.nvidia.com/books/HTML/gpugems/gpugems_ch38.html

A prova do teorema da decomposição de Helmholtz e uma enorme quantidade de material adicional sobre mecânica dos fluidos podem ser encontradas neste livro (em inglês, consulte a seção 1.2):
Chorin, AJ e JE Marsden. 1993. Uma Introdução Matemática à Mecânica dos Fluidos. 3rd ed. Springer

Canal de um YouTube de língua inglesa, criando conteúdo de qualidade relacionado à matemática e resolvendo equações diferenciais em particular (inglês). Vídeos muito visuais que ajudam a entender a essência de muitas coisas em matemática e física:
3Blue1Brown -
Equações diferenciais do YouTube (3Blue1Brown)

Também agradeço ao WhiteBlackGoose pela ajuda na preparação do material para o artigo.


E, no final, um pequeno bônus - algumas belas capturas de tela tiradas no programa:


Stream direto (configurações padrão)


Whirlpool (grande raio no applyForce)


Wave (alta vorticidade + difusão)

Além disso, pela demanda popular, adicionei um vídeo com a simulação:

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


All Articles