From 20eed2a258ece33f77c754750c3ae24eeb22157a Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 5 Aug 2023 18:34:48 +0900 Subject: [PATCH 01/13] add HLSL shaders --- .../Shaders/HLSL_DX12/bitonicSort.comp | 28 +++++ .../Fluid2D/Shaders/HLSL_DX12/buildGrid.comp | 45 ++++++++ .../Shaders/HLSL_DX12/buildGridIndices.comp | 30 +++++ .../Fluid2D/Shaders/HLSL_DX12/buildVBID.comp | 76 +++++++++++++ .../HLSL_DX12/calcCorrectPosition.comp | 99 ++++++++++++++++ .../Shaders/HLSL_DX12/calcExternalForce.comp | 30 +++++ .../Shaders/HLSL_DX12/calcScalingFactor.comp | 107 ++++++++++++++++++ .../Shaders/HLSL_DX12/clearGridIndices.comp | 8 ++ .../Fluid2D/Shaders/HLSL_DX12/integrate.comp | 22 ++++ .../Fluid2D/Shaders/HLSL_DX12/render.frag | 16 +++ 10 files changed, 461 insertions(+) create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/bitonicSort.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/buildGrid.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/buildGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/buildVBID.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/calcCorrectPosition.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/calcExternalForce.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/calcScalingFactor.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/clearGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/integrate.comp create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/render.frag diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/bitonicSort.comp b/examples/Fluid2D/Shaders/HLSL_DX12/bitonicSort.comp new file mode 100644 index 0000000..721a508 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/bitonicSort.comp @@ -0,0 +1,28 @@ +cbuffer CB : register(b0) +{ + float Inc; + float Dir; +}; + +RWStructuredBuffer gridTable : register(u0); + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + int inc = (int)Inc; + int dir = (int)Dir; + int t = dtid.x; + int low = t & (inc - 1); + int i = (t << 1) - low; + bool reverse = ((dir & i) == 0); + + int2 x0 = gridTable[i]; + int2 x1 = gridTable[inc + i]; + + int2 auxa = x0; + int2 auxb = x1; + if ((x0.x < x1.x) ^ reverse) { x0 = auxb; x1 = auxa; } + // Store + gridTable[i] = x0; + gridTable[inc + i] = x1; +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/buildGrid.comp b/examples/Fluid2D/Shaders/HLSL_DX12/buildGrid.comp new file mode 100644 index 0000000..7d1296f --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/buildGrid.comp @@ -0,0 +1,45 @@ +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + bool IsFix; +}; + +#define EMPTY_CELL 0x7fffffff + +int GetHash(int2 gridPos, int num) +{ + return gridPos.x + gridPos.y * num; +} + +int2 GetGridPos(float2 pos, float2 gridSize) +{ + return pos / gridSize; +} + +cbuffer CB : register(b0) +{ + int2 GridNum; + int2 GridSize; + int ParticlesCount; +}; + +RWStructuredBuffer particles : register(u0); +RWStructuredBuffer gridTable : register(u1); + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + if (dtid.x < ParticlesCount) + { + gridTable[dtid.x].x = GetHash(GetGridPos(particles[dtid.x].Next, GridSize), GridNum.x); + gridTable[dtid.x].y = dtid.x; + } + else + { + gridTable[dtid.x].x = EMPTY_CELL; + gridTable[dtid.x].y = dtid.x; + } +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/buildGridIndices.comp b/examples/Fluid2D/Shaders/HLSL_DX12/buildGridIndices.comp new file mode 100644 index 0000000..6cee450 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/buildGridIndices.comp @@ -0,0 +1,30 @@ +cbuffer CB : register(b0) +{ + int ParticlesCount; +}; + +RWStructuredBuffer gridTable : register(u0); +RWStructuredBuffer gridIndicesTable : register(u1); + +int getValidId(int id) { + return fmod(id + ParticlesCount, ParticlesCount); +} + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + int id = dtid.x; + int idPrev = getValidId(id - 1); + int idNext = getValidId(id + 1); + int cell = gridTable[id].x; + int cellPrev = gridTable[idPrev].x; + int cellNext = gridTable[idNext].x; + if (cell != cellPrev) + { + gridIndicesTable[cell].x = id; + } + if (cell != cellNext) + { + gridIndicesTable[cell].y = id + 1; + } +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/buildVBID.comp b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBID.comp new file mode 100644 index 0000000..578efe8 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBID.comp @@ -0,0 +1,76 @@ +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + bool IsFix; +}; + +#define EMPTY_CELL 0x7fffffff + +int GetHash(int2 gridPos, int num) +{ + return gridPos.x + gridPos.y * num; +} + +int2 GetGridPos(float2 pos, float2 gridSize) +{ + return pos / gridSize; +} + +struct Vertex +{ + float3 Position; + int Color; + float2 UV1; + float2 UV2; +}; + +cbuffer CB : register(b0) +{ + float ParticleRadius; + float4 Color; + float4 FixedColor; +}; + +RWStructuredBuffer particles : register(u0); +RWStructuredBuffer vertex_ : register(u1); +RWStructuredBuffer index : register(u2); + +int DecodeFloatRGBA(float4 rgba) { + int res = 0; + res += (int)(rgba.a * 255); + res = res << 8; + res += (int)(rgba.b * 255); + res = res << 8; + res += (int)(rgba.g * 255); + res = res << 8; + res += (int)(rgba.r * 255); + return res; +} + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + int c = DecodeFloatRGBA(lerp(Color, FixedColor, particles[dtid.x].IsFix)); + for (int i = 0; i < 4; i++) + { + vertex_[dtid.x * 4 + i].Color = c; + } + float3 pos = float3(particles[dtid.x].Current, 0.5); + vertex_[dtid.x * 4].Position = pos + float3(-1, -1, 0) * ParticleRadius * 3; + vertex_[dtid.x * 4 + 1].Position = pos + float3(1, -1, 0) * ParticleRadius * 3; + vertex_[dtid.x * 4 + 2].Position = pos + float3(1, 1, 0) * ParticleRadius * 3; + vertex_[dtid.x * 4 + 3].Position = pos + float3(-1, 1, 0) * ParticleRadius * 3; + vertex_[dtid.x * 4].UV1 = float2(0, 0); + vertex_[dtid.x * 4 + 1].UV1 = float2(1, 0); + vertex_[dtid.x * 4 + 2].UV1 = float2(1, 1); + vertex_[dtid.x * 4 + 3].UV1 = float2(0, 1); + index[dtid.x * 6] = dtid.x * 4; + index[dtid.x * 6 + 1] = dtid.x * 4 + 1; + index[dtid.x * 6 + 2] = dtid.x * 4 + 2; + index[dtid.x * 6 + 3] = dtid.x * 4; + index[dtid.x * 6 + 4] = dtid.x * 4 + 2; + index[dtid.x * 6 + 5] = dtid.x * 4 + 3; +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/calcCorrectPosition.comp b/examples/Fluid2D/Shaders/HLSL_DX12/calcCorrectPosition.comp new file mode 100644 index 0000000..8ebae5d --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/calcCorrectPosition.comp @@ -0,0 +1,99 @@ +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + bool IsFix; +}; + +#define EMPTY_CELL 0x7fffffff + +int GetHash(int2 gridPos, int num) +{ + return gridPos.x + gridPos.y * num; +} + +int2 GetGridPos(float2 pos, float2 gridSize) +{ + return pos / gridSize; +} + +cbuffer CB : register(b0) +{ + int2 GridNum; + int2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +}; + +RWStructuredBuffer particles : register(u0); +RWStructuredBuffer gridTable : register(u1); +RWStructuredBuffer gridIndicesTable : register(u2); + +float2 CalcPositionCorrectionCell(int2 gridPos, int i, float2 pos0) +{ + int gridHash = GetHash(gridPos, GridNum.x); + int startIndex = gridIndicesTable[gridHash].x; + float h = EffectiveRadius; + float r0 = Density; + float2 dp = float2(0.0f, 0.0f); + float dt = Dt; + float si = particles[i].Pscl; + if(startIndex != EMPTY_CELL){ // セルが空でないかのチェック + // セル内のパーティクルで反復 + uint endIndex = gridIndicesTable[gridHash].y; + for(uint j = startIndex; j < endIndex; ++j){ + if(gridTable[j].y == i) continue; + float2 pos1 = particles[gridTable[j].y].Next; + float2 rij = pos0 - pos1; + float r = length(rij); + if(r <= h && r > 0.0){ + float scorr = 0; + { + float q = h * h - r * r; + float q2 = h * h - 0.04 * h * h; + float ww = Wpoly6 * q * q * q / (Wpoly6 * q2 * q2 * q2); + scorr = -0.1 * pow(ww, 4) * dt * dt; + } + { + float q = h - r; + float sj = particles[gridTable[j].y].Pscl; + // Spikyカーネルで位置修正量を計算 + dp += (si + sj + scorr) * (GWspiky * q * q * rij / r) / r0; + } + } + } + } + return dp; +} + +float2 CalcPositionCorrection(int id) +{ + float2 pos = particles[id].Next; // パーティクル位置 + float h = EffectiveRadius; + // パーティクル周囲のグリッド + int2 grid_pos0 = GetGridPos(pos - h, GridSize.xy); + int2 grid_pos1 = GetGridPos(pos + h, GridSize.xy); + // 周囲のグリッドも含めて近傍探索，位置修正量を計算 + float2 dpij = float2(0.0f, 0.0f); + for(int y = grid_pos0.y; y <= grid_pos1.y; ++y){ + for(int x = grid_pos0.x; x <= grid_pos1.x; ++x){ + int2 n_grid_pos = int2(x, y); + dpij += CalcPositionCorrectionCell(n_grid_pos, id, pos); + } + } + return dpij; +} + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + int id = dtid.x; + if (particles[id].IsFix) return; + particles[id].Next += CalcPositionCorrection(id); +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/calcExternalForce.comp b/examples/Fluid2D/Shaders/HLSL_DX12/calcExternalForce.comp new file mode 100644 index 0000000..83cb76b --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/calcExternalForce.comp @@ -0,0 +1,30 @@ +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + bool IsFix; +}; + +cbuffer CB : register(b0) +{ + float2 Force; + float2 Gravity; + float2 Dt; + int2 GridNum; + int2 GridSize; +}; + +RWStructuredBuffer particles : register(u0); + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + if (particles[dtid.x].IsFix) return; + particles[dtid.x].Velocity += (Gravity + Force) * Dt.x; + float2 pos = particles[dtid.x].Current + particles[dtid.x].Velocity * Dt.x; + pos.x = clamp(pos.x, 8, GridNum.x * GridSize.x - 8); + pos.y = clamp(pos.y, 8, GridNum.y * GridSize.y - 8); + particles[dtid.x].Next = pos; +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/calcScalingFactor.comp b/examples/Fluid2D/Shaders/HLSL_DX12/calcScalingFactor.comp new file mode 100644 index 0000000..edb7d70 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/calcScalingFactor.comp @@ -0,0 +1,107 @@ +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + bool IsFix; +}; + +#define EMPTY_CELL 0x7fffffff + +int GetHash(int2 gridPos, int num) +{ + return gridPos.x + gridPos.y * num; +} + +int2 GetGridPos(float2 pos, float2 gridSize) +{ + return pos / gridSize; +} + +cbuffer CB : register(b0) +{ + int2 GridNum; + int2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +}; +RWStructuredBuffer particles : register(u0); +RWStructuredBuffer gridTable : register(u1); +RWStructuredBuffer gridIndicesTable : register(u2); + +float2 CalcDensityCellPB(int2 gridPos, int i, float2 pos0) +{ + int gridHash = GetHash(gridPos, GridNum.x); + int startIndex = gridIndicesTable[gridHash].x; + float h2 = EffectiveRadius * EffectiveRadius; + float dens = 0.0f; + if(startIndex != EMPTY_CELL){ + int endIndex = gridIndicesTable[gridHash].y; + for(int j = startIndex; j < endIndex; ++j){ + float2 pos1 = particles[gridTable[j].y].Next; + float r2 = dot(pos0 - pos1, pos0 - pos1); + if(r2 <= h2){ + float q = h2 - r2; + dens += Wpoly6 * q * q * q; + } + } + } + return dens; +} + +float CalcScalingFactorCell(int2 gridPos, int i, float2 pos0) +{ + int gridHash = GetHash(gridPos, GridNum.x); + int startIndex = gridIndicesTable[gridHash].x; + float h = EffectiveRadius; + float r0 = Density; + float sd = 0.0f; + float2 sd2 = float2(0.0f, 0.0f); + if(startIndex != EMPTY_CELL){ + uint endIndex = gridIndicesTable[gridHash].y; + for(uint j = startIndex; j < endIndex; ++j){ + float2 pos1 = particles[gridTable[j].y].Next; + float2 rij = pos0 - pos1; + float r = length(rij); + if(r <= h && r > 0.0){ + float q = h - r; + float2 dp = (GWspiky * q * q * rij / r) / r0; + sd2 += dp; + sd += dot(dp, dp); + } + } + } + return sd + dot(sd2, sd2); +} + +void CalcScalingFactor(int id) +{ + float2 pos = particles[id].Next; + float h = EffectiveRadius; + float r0 = Density; + int2 grid_pos0 = GetGridPos(pos-float2(h, h), GridSize); + int2 grid_pos1 = GetGridPos(pos+float2(h, h), GridSize); + + float dens = 0.0f; + float sd = 0.0f; + for(int y = grid_pos0.y; y <= grid_pos1.y; ++y){ + for(int x = grid_pos0.x; x <= grid_pos1.x; ++x){ + int2 n_grid_pos = int2(x, y); + dens += CalcDensityCellPB(n_grid_pos, id, pos); + sd += CalcScalingFactorCell(n_grid_pos, id, pos); + } + } + float C = dens/r0-1.0; + particles[id].Pscl = -C/(sd+Eps); +} +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + int id = dtid.x; + CalcScalingFactor(id); +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/clearGridIndices.comp b/examples/Fluid2D/Shaders/HLSL_DX12/clearGridIndices.comp new file mode 100644 index 0000000..fb535de --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/clearGridIndices.comp @@ -0,0 +1,8 @@ +#define EMPTY_CELL 0x7fffffff + +RWStructuredBuffer gridIndicesTable : register(u0); +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + gridIndicesTable[dtid.x] = int2(EMPTY_CELL, EMPTY_CELL); +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/integrate.comp b/examples/Fluid2D/Shaders/HLSL_DX12/integrate.comp new file mode 100644 index 0000000..f90cb75 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/integrate.comp @@ -0,0 +1,22 @@ +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + bool IsFix; +}; + +cbuffer CB : register(b0) +{ + float Dt; +}; + +RWStructuredBuffer particles : register(u0); + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + particles[dtid.x].Velocity = (particles[dtid.x].Next - particles[dtid.x].Current) / Dt; + particles[dtid.x].Current = particles[dtid.x].Next; +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/render.frag b/examples/Fluid2D/Shaders/HLSL_DX12/render.frag new file mode 100644 index 0000000..9e41453 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/render.frag @@ -0,0 +1,16 @@ +struct PS_INPUT +{ + float4 Position : SV_POSITION; + float4 Color : COLOR0; + float2 UV1 : UV0; + float2 UV2 : UV1; +}; + +float4 main(PS_INPUT input) : SV_TARGET +{ + float4 c; + float r = length((input.UV1 - 0.5) * 2); + float a = r > 1 ? 0 : (-4 / 9 * pow(r, 6) + 17 / 9 * pow(r, 4) - 22 / 9 * pow(r, 2) + 1); + c = float4(input.Color.rgb, input.Color.a * a); + return c; +} \ No newline at end of file From 0ae09ce4a2acc15d2ce4a1fae6a7f446af679a1a Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 5 Aug 2023 18:35:14 +0900 Subject: [PATCH 02/13] transpile shaders for other platforms --- .../Fluid2D/Shaders/GLSL_GL/bitonicSort.comp | 42 +++++ .../Fluid2D/Shaders/GLSL_GL/buildGrid.comp | 64 +++++++ .../Shaders/GLSL_GL/buildGridIndices.comp | 50 ++++++ .../Fluid2D/Shaders/GLSL_GL/buildVBID.comp | 87 ++++++++++ .../Shaders/GLSL_GL/calcCorrectPosition.comp | 132 +++++++++++++++ .../Shaders/GLSL_GL/calcExternalForce.comp | 46 ++++++ .../Shaders/GLSL_GL/calcScalingFactor.comp | 151 +++++++++++++++++ .../Shaders/GLSL_GL/clearGridIndices.comp | 20 +++ .../Fluid2D/Shaders/GLSL_GL/integrate.comp | 35 ++++ examples/Fluid2D/Shaders/GLSL_GL/render.frag | 34 ++++ .../Shaders/GLSL_VULKAN/bitonicSort.comp | 42 +++++ .../Shaders/GLSL_VULKAN/buildGrid.comp | 64 +++++++ .../Shaders/GLSL_VULKAN/buildGridIndices.comp | 50 ++++++ .../Shaders/GLSL_VULKAN/buildVBID.comp | 87 ++++++++++ .../GLSL_VULKAN/calcCorrectPosition.comp | 132 +++++++++++++++ .../GLSL_VULKAN/calcExternalForce.comp | 46 ++++++ .../GLSL_VULKAN/calcScalingFactor.comp | 151 +++++++++++++++++ .../Shaders/GLSL_VULKAN/clearGridIndices.comp | 20 +++ .../Shaders/GLSL_VULKAN/integrate.comp | 35 ++++ .../Fluid2D/Shaders/GLSL_VULKAN/render.frag | 34 ++++ .../Fluid2D/Shaders/Metal/bitonicSort.comp | 47 ++++++ examples/Fluid2D/Shaders/Metal/buildGrid.comp | 71 ++++++++ .../Shaders/Metal/buildGridIndices.comp | 58 +++++++ examples/Fluid2D/Shaders/Metal/buildVBID.comp | 93 +++++++++++ .../Shaders/Metal/calcCorrectPosition.comp | 136 +++++++++++++++ .../Shaders/Metal/calcExternalForce.comp | 51 ++++++ .../Shaders/Metal/calcScalingFactor.comp | 156 ++++++++++++++++++ .../Shaders/Metal/clearGridIndices.comp | 25 +++ examples/Fluid2D/Shaders/Metal/integrate.comp | 40 +++++ examples/Fluid2D/Shaders/Metal/render.frag | 49 ++++++ .../Shaders/SPIRV/bitonicSort.comp.spv | Bin 0 -> 2700 bytes .../Fluid2D/Shaders/SPIRV/buildGrid.comp.spv | Bin 0 -> 3532 bytes .../Shaders/SPIRV/buildGridIndices.comp.spv | Bin 0 -> 2868 bytes .../Fluid2D/Shaders/SPIRV/buildVBID.comp.spv | Bin 0 -> 6956 bytes .../SPIRV/calcCorrectPosition.comp.spv | Bin 0 -> 9456 bytes .../Shaders/SPIRV/calcExternalForce.comp.spv | Bin 0 -> 3048 bytes .../Shaders/SPIRV/calcScalingFactor.comp.spv | Bin 0 -> 10812 bytes .../Shaders/SPIRV/clearGridIndices.comp.spv | Bin 0 -> 1096 bytes .../Fluid2D/Shaders/SPIRV/integrate.comp.spv | Bin 0 -> 1920 bytes .../Fluid2D/Shaders/SPIRV/render.frag.spv | Bin 0 -> 2388 bytes 40 files changed, 2048 insertions(+) create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/bitonicSort.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/buildGrid.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/buildGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/buildVBID.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/calcCorrectPosition.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/calcExternalForce.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/calcScalingFactor.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/clearGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/integrate.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/render.frag create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBID.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/render.frag create mode 100644 examples/Fluid2D/Shaders/Metal/bitonicSort.comp create mode 100644 examples/Fluid2D/Shaders/Metal/buildGrid.comp create mode 100644 examples/Fluid2D/Shaders/Metal/buildGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/Metal/buildVBID.comp create mode 100644 examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp create mode 100644 examples/Fluid2D/Shaders/Metal/calcExternalForce.comp create mode 100644 examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp create mode 100644 examples/Fluid2D/Shaders/Metal/clearGridIndices.comp create mode 100644 examples/Fluid2D/Shaders/Metal/integrate.comp create mode 100644 examples/Fluid2D/Shaders/Metal/render.frag create mode 100644 examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/buildGrid.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/buildGridIndices.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/buildVBID.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/calcCorrectPosition.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/calcExternalForce.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/calcScalingFactor.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/clearGridIndices.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv create mode 100644 examples/Fluid2D/Shaders/SPIRV/render.frag.spv diff --git a/examples/Fluid2D/Shaders/GLSL_GL/bitonicSort.comp b/examples/Fluid2D/Shaders/GLSL_GL/bitonicSort.comp new file mode 100644 index 0000000..b0f6758 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/bitonicSort.comp @@ -0,0 +1,42 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 0, std140) uniform CB +{ + float Inc; + float Dir; +} CBCS0; + +layout(binding = 0, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +void _main(uvec3 dtid) +{ + int inc = int(CBCS0.Inc); + int dir = int(CBCS0.Dir); + int t = int(dtid.x); + int low = t & (inc - 1); + int i = (t << 1) - low; + bool reverse = (dir & i) == 0; + ivec2 x0 = gridTable_1._data[i]; + ivec2 x1 = gridTable_1._data[inc + i]; + ivec2 auxa = x0; + ivec2 auxb = x1; + if ((int(x0.x < x1.x) ^ int(reverse)) != int(0u)) + { + x0 = auxb; + x1 = auxa; + } + gridTable_1._data[i] = x0; + gridTable_1._data[inc + i] = x1; +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/buildGrid.comp b/examples/Fluid2D/Shaders/GLSL_GL/buildGrid.comp new file mode 100644 index 0000000..d350c44 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/buildGrid.comp @@ -0,0 +1,64 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(binding = 0, std140) uniform CB +{ + ivec2 GridNum; + ivec2 GridSize; + int ParticlesCount; +} CBCS0; + +layout(binding = 1, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +layout(binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +ivec2 GetGridPos(vec2 pos, vec2 gridSize) +{ + return ivec2(pos / gridSize); +} + +int GetHash(ivec2 gridPos, int num) +{ + return gridPos.x + (gridPos.y * num); +} + +void _main(uvec3 dtid) +{ + if (dtid.x < uint(CBCS0.ParticlesCount)) + { + vec2 param = particles_1._data[dtid.x].Next; + vec2 param_1 = vec2(CBCS0.GridSize); + ivec2 param_2 = GetGridPos(param, param_1); + int param_3 = CBCS0.GridNum.x; + gridTable_1._data[dtid.x].x = GetHash(param_2, param_3); + gridTable_1._data[dtid.x].y = int(dtid.x); + } + else + { + gridTable_1._data[dtid.x].x = 2147483647; + gridTable_1._data[dtid.x].y = int(dtid.x); + } +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/buildGridIndices.comp b/examples/Fluid2D/Shaders/GLSL_GL/buildGridIndices.comp new file mode 100644 index 0000000..37350f9 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/buildGridIndices.comp @@ -0,0 +1,50 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 0, std140) uniform CB +{ + int ParticlesCount; +} CBCS0; + +layout(binding = 0, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +layout(binding = 1, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +int getValidId(int id) +{ + return int(mod(float(id + CBCS0.ParticlesCount), float(CBCS0.ParticlesCount))); +} + +void _main(uvec3 dtid) +{ + int id = int(dtid.x); + int param = id - 1; + int idPrev = getValidId(param); + int param_1 = id + 1; + int idNext = getValidId(param_1); + int cell = gridTable_1._data[id].x; + int cellPrev = gridTable_1._data[idPrev].x; + int cellNext = gridTable_1._data[idNext].x; + if (cell != cellPrev) + { + gridIndicesTable_1._data[cell].x = id; + } + if (cell != cellNext) + { + gridIndicesTable_1._data[cell].y = id + 1; + } +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/buildVBID.comp b/examples/Fluid2D/Shaders/GLSL_GL/buildVBID.comp new file mode 100644 index 0000000..79977c1 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/buildVBID.comp @@ -0,0 +1,87 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +struct Vertex +{ + vec3 Position; + int Color; + vec2 UV1; + vec2 UV2; +}; + +layout(binding = 0, std140) uniform CB +{ + float ParticleRadius; + vec4 Color; + vec4 FixedColor; +} CBCS0; + +layout(binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(binding = 1, std430) buffer vertex_ +{ + Vertex _data[]; +} vertex_1; + +layout(binding = 2, std430) buffer index +{ + int _data[]; +} index_1; + +int DecodeFloatRGBA(vec4 rgba) +{ + int res = 0; + res += int(rgba.w * 255.0); + res = res << 8; + res += int(rgba.z * 255.0); + res = res << 8; + res += int(rgba.y * 255.0); + res = res << 8; + res += int(rgba.x * 255.0); + return res; +} + +void _main(uvec3 dtid) +{ + vec4 param = mix(CBCS0.Color, CBCS0.FixedColor, vec4(float(particles_1._data[dtid.x].IsFix != 0u))); + int c = DecodeFloatRGBA(param); + for (int i = 0; i < 4; i++) + { + vertex_1._data[(dtid.x * 4u) + uint(i)].Color = c; + } + vec3 pos = vec3(particles_1._data[dtid.x].Current, 0.5); + vertex_1._data[dtid.x * 4u].Position = pos + ((vec3(-1.0, -1.0, 0.0) * CBCS0.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 1u].Position = pos + ((vec3(1.0, -1.0, 0.0) * CBCS0.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 2u].Position = pos + ((vec3(1.0, 1.0, 0.0) * CBCS0.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 3u].Position = pos + ((vec3(-1.0, 1.0, 0.0) * CBCS0.ParticleRadius) * 3.0); + vertex_1._data[dtid.x * 4u].UV1 = vec2(0.0); + vertex_1._data[(dtid.x * 4u) + 1u].UV1 = vec2(1.0, 0.0); + vertex_1._data[(dtid.x * 4u) + 2u].UV1 = vec2(1.0); + vertex_1._data[(dtid.x * 4u) + 3u].UV1 = vec2(0.0, 1.0); + index_1._data[dtid.x * 6u] = int(dtid.x * 4u); + index_1._data[(dtid.x * 6u) + 1u] = int((dtid.x * 4u) + 1u); + index_1._data[(dtid.x * 6u) + 2u] = int((dtid.x * 4u) + 2u); + index_1._data[(dtid.x * 6u) + 3u] = int(dtid.x * 4u); + index_1._data[(dtid.x * 6u) + 4u] = int((dtid.x * 4u) + 2u); + index_1._data[(dtid.x * 6u) + 5u] = int((dtid.x * 4u) + 3u); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/calcCorrectPosition.comp b/examples/Fluid2D/Shaders/GLSL_GL/calcCorrectPosition.comp new file mode 100644 index 0000000..850a5a9 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/calcCorrectPosition.comp @@ -0,0 +1,132 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(binding = 0, std140) uniform CB +{ + ivec2 GridNum; + ivec2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +} CBCS0; + +layout(binding = 2, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +layout(binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(binding = 1, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +ivec2 GetGridPos(vec2 pos, vec2 gridSize) +{ + return ivec2(pos / gridSize); +} + +int GetHash(ivec2 gridPos, int num) +{ + return gridPos.x + (gridPos.y * num); +} + +vec2 CalcPositionCorrectionCell(ivec2 gridPos, int i, vec2 pos0) +{ + ivec2 param = gridPos; + int param_1 = CBCS0.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = gridIndicesTable_1._data[gridHash].x; + float h = CBCS0.EffectiveRadius; + float r0 = CBCS0.Density; + vec2 dp = vec2(0.0); + float dt = CBCS0.Dt; + float si = particles_1._data[i].Pscl; + if (startIndex != 2147483647) + { + uint endIndex = uint(gridIndicesTable_1._data[gridHash].y); + uint _117 = uint(startIndex); + for (uint j = _117; j < endIndex; j++) + { + if (gridTable_1._data[j].y == i) + { + continue; + } + vec2 pos1 = particles_1._data[gridTable_1._data[j].y].Next; + vec2 rij = pos0 - pos1; + float r = length(rij); + if ((r <= h) && (r > 0.0)) + { + float scorr = 0.0; + float q = (h * h) - (r * r); + float q2 = (h * h) - ((0.039999999105930328369140625 * h) * h); + float ww = (((CBCS0.Wpoly6 * q) * q) * q) / (((CBCS0.Wpoly6 * q2) * q2) * q2); + scorr = (((-0.100000001490116119384765625) * pow(ww, 4.0)) * dt) * dt; + float q_1 = h - r; + float sj = particles_1._data[gridTable_1._data[j].y].Pscl; + dp += ((((rij * ((CBCS0.GWspiky * q_1) * q_1)) / vec2(r)) * ((si + sj) + scorr)) / vec2(r0)); + } + } + } + return dp; +} + +vec2 CalcPositionCorrection(int id) +{ + vec2 pos = particles_1._data[id].Next; + float h = CBCS0.EffectiveRadius; + vec2 param = pos - vec2(h); + vec2 param_1 = vec2(CBCS0.GridSize); + ivec2 grid_pos0 = GetGridPos(param, param_1); + vec2 param_2 = pos + vec2(h); + vec2 param_3 = vec2(CBCS0.GridSize); + ivec2 grid_pos1 = GetGridPos(param_2, param_3); + vec2 dpij = vec2(0.0); + for (int y = grid_pos0.y; y <= grid_pos1.y; y++) + { + for (int x = grid_pos0.x; x <= grid_pos1.x; x++) + { + ivec2 n_grid_pos = ivec2(x, y); + ivec2 param_4 = n_grid_pos; + int param_5 = id; + vec2 param_6 = pos; + dpij += CalcPositionCorrectionCell(param_4, param_5, param_6); + } + } + return dpij; +} + +void _main(uvec3 dtid) +{ + int id = int(dtid.x); + if (particles_1._data[id].IsFix != 0u) + { + return; + } + int param = id; + particles_1._data[id].Next += CalcPositionCorrection(param); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/calcExternalForce.comp b/examples/Fluid2D/Shaders/GLSL_GL/calcExternalForce.comp new file mode 100644 index 0000000..e377d5a --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/calcExternalForce.comp @@ -0,0 +1,46 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(binding = 0, std140) uniform CB +{ + vec2 Force; + vec2 Gravity; + vec2 Dt; + ivec2 GridNum; + ivec2 GridSize; +} CBCS0; + +void _main(uvec3 dtid) +{ + if (particles_1._data[dtid.x].IsFix != 0u) + { + return; + } + particles_1._data[dtid.x].Velocity += ((CBCS0.Gravity + CBCS0.Force) * CBCS0.Dt.x); + vec2 pos = particles_1._data[dtid.x].Current + (particles_1._data[dtid.x].Velocity * CBCS0.Dt.x); + pos.x = clamp(pos.x, 8.0, float((CBCS0.GridNum.x * CBCS0.GridSize.x) - 8)); + pos.y = clamp(pos.y, 8.0, float((CBCS0.GridNum.y * CBCS0.GridSize.y) - 8)); + particles_1._data[dtid.x].Next = pos; +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/calcScalingFactor.comp b/examples/Fluid2D/Shaders/GLSL_GL/calcScalingFactor.comp new file mode 100644 index 0000000..42e1da2 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/calcScalingFactor.comp @@ -0,0 +1,151 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(binding = 0, std140) uniform CB +{ + ivec2 GridNum; + ivec2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +} CBCS0; + +layout(binding = 2, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +layout(binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(binding = 1, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +ivec2 GetGridPos(vec2 pos, vec2 gridSize) +{ + return ivec2(pos / gridSize); +} + +int GetHash(ivec2 gridPos, int num) +{ + return gridPos.x + (gridPos.y * num); +} + +vec2 CalcDensityCellPB(ivec2 gridPos, int i, vec2 pos0) +{ + ivec2 param = gridPos; + int param_1 = CBCS0.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = gridIndicesTable_1._data[gridHash].x; + float h2 = CBCS0.EffectiveRadius * CBCS0.EffectiveRadius; + float dens = 0.0; + if (startIndex != 2147483647) + { + int endIndex = gridIndicesTable_1._data[gridHash].y; + for (int j = startIndex; j < endIndex; j++) + { + vec2 pos1 = particles_1._data[gridTable_1._data[j].y].Next; + float r2 = dot(pos0 - pos1, pos0 - pos1); + if (r2 <= h2) + { + float q = h2 - r2; + dens += (((CBCS0.Wpoly6 * q) * q) * q); + } + } + } + return vec2(dens); +} + +float CalcScalingFactorCell(ivec2 gridPos, int i, vec2 pos0) +{ + ivec2 param = gridPos; + int param_1 = CBCS0.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = gridIndicesTable_1._data[gridHash].x; + float h = CBCS0.EffectiveRadius; + float r0 = CBCS0.Density; + float sd = 0.0; + vec2 sd2 = vec2(0.0); + if (startIndex != 2147483647) + { + uint endIndex = uint(gridIndicesTable_1._data[gridHash].y); + uint _195 = uint(startIndex); + for (uint j = _195; j < endIndex; j++) + { + vec2 pos1 = particles_1._data[gridTable_1._data[j].y].Next; + vec2 rij = pos0 - pos1; + float r = length(rij); + if ((r <= h) && (r > 0.0)) + { + float q = h - r; + vec2 dp = ((rij * ((CBCS0.GWspiky * q) * q)) / vec2(r)) / vec2(r0); + sd2 += dp; + sd += dot(dp, dp); + } + } + } + return sd + dot(sd2, sd2); +} + +void CalcScalingFactor(int id) +{ + vec2 pos = particles_1._data[id].Next; + float h = CBCS0.EffectiveRadius; + float r0 = CBCS0.Density; + vec2 param = pos - vec2(h, h); + vec2 param_1 = vec2(CBCS0.GridSize); + ivec2 grid_pos0 = GetGridPos(param, param_1); + vec2 param_2 = pos + vec2(h, h); + vec2 param_3 = vec2(CBCS0.GridSize); + ivec2 grid_pos1 = GetGridPos(param_2, param_3); + float dens = 0.0; + float sd = 0.0; + for (int y = grid_pos0.y; y <= grid_pos1.y; y++) + { + for (int x = grid_pos0.x; x <= grid_pos1.x; x++) + { + ivec2 n_grid_pos = ivec2(x, y); + ivec2 param_4 = n_grid_pos; + int param_5 = id; + vec2 param_6 = pos; + dens += CalcDensityCellPB(param_4, param_5, param_6).x; + ivec2 param_7 = n_grid_pos; + int param_8 = id; + vec2 param_9 = pos; + sd += CalcScalingFactorCell(param_7, param_8, param_9); + } + } + float C = (dens / r0) - 1.0; + particles_1._data[id].Pscl = (-C) / (sd + CBCS0.Eps); +} + +void _main(uvec3 dtid) +{ + int id = int(dtid.x); + int param = id; + CalcScalingFactor(param); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/clearGridIndices.comp b/examples/Fluid2D/Shaders/GLSL_GL/clearGridIndices.comp new file mode 100644 index 0000000..2b36c3a --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/clearGridIndices.comp @@ -0,0 +1,20 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 0, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +void _main(uvec3 dtid) +{ + gridIndicesTable_1._data[dtid.x] = ivec2(2147483647); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/integrate.comp b/examples/Fluid2D/Shaders/GLSL_GL/integrate.comp new file mode 100644 index 0000000..22db3bd --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/integrate.comp @@ -0,0 +1,35 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(binding = 0, std140) uniform CB +{ + float Dt; +} CBCS0; + +void _main(uvec3 dtid) +{ + particles_1._data[dtid.x].Velocity = (particles_1._data[dtid.x].Next - particles_1._data[dtid.x].Current) / vec2(CBCS0.Dt); + particles_1._data[dtid.x].Current = particles_1._data[dtid.x].Next; +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_GL/render.frag b/examples/Fluid2D/Shaders/GLSL_GL/render.frag new file mode 100644 index 0000000..aa0abac --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/render.frag @@ -0,0 +1,34 @@ +#version 430 + +struct PS_INPUT +{ + vec4 Position; + vec4 Color; + vec2 UV1; + vec2 UV2; +}; + +layout(location = 0) in vec4 input_Color; +layout(location = 1) in vec2 input_UV1; +layout(location = 2) in vec2 input_UV2; +layout(location = 0) out vec4 _entryPointOutput; + +vec4 _main(PS_INPUT _input) +{ + float r = length((_input.UV1 - vec2(0.5)) * 2.0); + float a = (r > 1.0) ? 0.0 : ((((0.0 * pow(r, 6.0)) + (1.0 * pow(r, 4.0))) - (2.0 * pow(r, 2.0))) + 1.0); + vec4 c = vec4(_input.Color.xyz, _input.Color.w * a); + return c; +} + +void main() +{ + PS_INPUT _input; + _input.Position = gl_FragCoord; + _input.Color = input_Color; + _input.UV1 = input_UV1; + _input.UV2 = input_UV2; + PS_INPUT param = _input; + _entryPointOutput = _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp new file mode 100644 index 0000000..3eef971 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp @@ -0,0 +1,42 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(set = 0, binding = 0, std140) uniform CB +{ + float Inc; + float Dir; +} _19; + +layout(set = 0, binding = 1, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +void _main(uvec3 dtid) +{ + int inc = int(_19.Inc); + int dir = int(_19.Dir); + int t = int(dtid.x); + int low = t & (inc - 1); + int i = (t << 1) - low; + bool reverse = (dir & i) == 0; + ivec2 x0 = gridTable_1._data[i]; + ivec2 x1 = gridTable_1._data[inc + i]; + ivec2 auxa = x0; + ivec2 auxb = x1; + if ((int(x0.x < x1.x) ^ int(reverse)) != int(0u)) + { + x0 = auxb; + x1 = auxa; + } + gridTable_1._data[i] = x0; + gridTable_1._data[inc + i] = x1; +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp new file mode 100644 index 0000000..7964108 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp @@ -0,0 +1,64 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(set = 0, binding = 0, std140) uniform CB +{ + ivec2 GridNum; + ivec2 GridSize; + int ParticlesCount; +} _52; + +layout(set = 0, binding = 2, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +layout(set = 0, binding = 1, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +ivec2 GetGridPos(vec2 pos, vec2 gridSize) +{ + return ivec2(pos / gridSize); +} + +int GetHash(ivec2 gridPos, int num) +{ + return gridPos.x + (gridPos.y * num); +} + +void _main(uvec3 dtid) +{ + if (dtid.x < uint(_52.ParticlesCount)) + { + vec2 param = particles_1._data[dtid.x].Next; + vec2 param_1 = vec2(_52.GridSize); + ivec2 param_2 = GetGridPos(param, param_1); + int param_3 = _52.GridNum.x; + gridTable_1._data[dtid.x].x = GetHash(param_2, param_3); + gridTable_1._data[dtid.x].y = int(dtid.x); + } + else + { + gridTable_1._data[dtid.x].x = 2147483647; + gridTable_1._data[dtid.x].y = int(dtid.x); + } +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp new file mode 100644 index 0000000..8920fad --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp @@ -0,0 +1,50 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(set = 0, binding = 0, std140) uniform CB +{ + int ParticlesCount; +} _22; + +layout(set = 0, binding = 1, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +layout(set = 0, binding = 2, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +int getValidId(int id) +{ + return int(mod(float(id + _22.ParticlesCount), float(_22.ParticlesCount))); +} + +void _main(uvec3 dtid) +{ + int id = int(dtid.x); + int param = id - 1; + int idPrev = getValidId(param); + int param_1 = id + 1; + int idNext = getValidId(param_1); + int cell = gridTable_1._data[id].x; + int cellPrev = gridTable_1._data[idPrev].x; + int cellNext = gridTable_1._data[idNext].x; + if (cell != cellPrev) + { + gridIndicesTable_1._data[cell].x = id; + } + if (cell != cellNext) + { + gridIndicesTable_1._data[cell].y = id + 1; + } +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBID.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBID.comp new file mode 100644 index 0000000..32d0d47 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBID.comp @@ -0,0 +1,87 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +struct Vertex +{ + vec3 Position; + int Color; + vec2 UV1; + vec2 UV2; +}; + +layout(set = 0, binding = 0, std140) uniform CB +{ + float ParticleRadius; + vec4 Color; + vec4 FixedColor; +} _67; + +layout(set = 0, binding = 1, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 0, binding = 2, std430) buffer vertex_ +{ + Vertex _data[]; +} vertex_1; + +layout(set = 0, binding = 3, std430) buffer index +{ + int _data[]; +} index_1; + +int DecodeFloatRGBA(vec4 rgba) +{ + int res = 0; + res += int(rgba.w * 255.0); + res = res << 8; + res += int(rgba.z * 255.0); + res = res << 8; + res += int(rgba.y * 255.0); + res = res << 8; + res += int(rgba.x * 255.0); + return res; +} + +void _main(uvec3 dtid) +{ + vec4 param = mix(_67.Color, _67.FixedColor, vec4(float(particles_1._data[dtid.x].IsFix != 0u))); + int c = DecodeFloatRGBA(param); + for (int i = 0; i < 4; i++) + { + vertex_1._data[(dtid.x * 4u) + uint(i)].Color = c; + } + vec3 pos = vec3(particles_1._data[dtid.x].Current, 0.5); + vertex_1._data[dtid.x * 4u].Position = pos + ((vec3(-1.0, -1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 1u].Position = pos + ((vec3(1.0, -1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 2u].Position = pos + ((vec3(1.0, 1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 3u].Position = pos + ((vec3(-1.0, 1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[dtid.x * 4u].UV1 = vec2(0.0); + vertex_1._data[(dtid.x * 4u) + 1u].UV1 = vec2(1.0, 0.0); + vertex_1._data[(dtid.x * 4u) + 2u].UV1 = vec2(1.0); + vertex_1._data[(dtid.x * 4u) + 3u].UV1 = vec2(0.0, 1.0); + index_1._data[dtid.x * 6u] = int(dtid.x * 4u); + index_1._data[(dtid.x * 6u) + 1u] = int((dtid.x * 4u) + 1u); + index_1._data[(dtid.x * 6u) + 2u] = int((dtid.x * 4u) + 2u); + index_1._data[(dtid.x * 6u) + 3u] = int(dtid.x * 4u); + index_1._data[(dtid.x * 6u) + 4u] = int((dtid.x * 4u) + 2u); + index_1._data[(dtid.x * 6u) + 5u] = int((dtid.x * 4u) + 3u); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp new file mode 100644 index 0000000..257d97a --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp @@ -0,0 +1,132 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(set = 0, binding = 0, std140) uniform CB +{ + ivec2 GridNum; + ivec2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +} _60; + +layout(set = 0, binding = 3, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +layout(set = 0, binding = 1, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 0, binding = 2, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +ivec2 GetGridPos(vec2 pos, vec2 gridSize) +{ + return ivec2(pos / gridSize); +} + +int GetHash(ivec2 gridPos, int num) +{ + return gridPos.x + (gridPos.y * num); +} + +vec2 CalcPositionCorrectionCell(ivec2 gridPos, int i, vec2 pos0) +{ + ivec2 param = gridPos; + int param_1 = _60.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = gridIndicesTable_1._data[gridHash].x; + float h = _60.EffectiveRadius; + float r0 = _60.Density; + vec2 dp = vec2(0.0); + float dt = _60.Dt; + float si = particles_1._data[i].Pscl; + if (startIndex != 2147483647) + { + uint endIndex = uint(gridIndicesTable_1._data[gridHash].y); + uint _117 = uint(startIndex); + for (uint j = _117; j < endIndex; j++) + { + if (gridTable_1._data[j].y == i) + { + continue; + } + vec2 pos1 = particles_1._data[gridTable_1._data[j].y].Next; + vec2 rij = pos0 - pos1; + float r = length(rij); + if ((r <= h) && (r > 0.0)) + { + float scorr = 0.0; + float q = (h * h) - (r * r); + float q2 = (h * h) - ((0.039999999105930328369140625 * h) * h); + float ww = (((_60.Wpoly6 * q) * q) * q) / (((_60.Wpoly6 * q2) * q2) * q2); + scorr = (((-0.100000001490116119384765625) * pow(ww, 4.0)) * dt) * dt; + float q_1 = h - r; + float sj = particles_1._data[gridTable_1._data[j].y].Pscl; + dp += ((((rij * ((_60.GWspiky * q_1) * q_1)) / vec2(r)) * ((si + sj) + scorr)) / vec2(r0)); + } + } + } + return dp; +} + +vec2 CalcPositionCorrection(int id) +{ + vec2 pos = particles_1._data[id].Next; + float h = _60.EffectiveRadius; + vec2 param = pos - vec2(h); + vec2 param_1 = vec2(_60.GridSize); + ivec2 grid_pos0 = GetGridPos(param, param_1); + vec2 param_2 = pos + vec2(h); + vec2 param_3 = vec2(_60.GridSize); + ivec2 grid_pos1 = GetGridPos(param_2, param_3); + vec2 dpij = vec2(0.0); + for (int y = grid_pos0.y; y <= grid_pos1.y; y++) + { + for (int x = grid_pos0.x; x <= grid_pos1.x; x++) + { + ivec2 n_grid_pos = ivec2(x, y); + ivec2 param_4 = n_grid_pos; + int param_5 = id; + vec2 param_6 = pos; + dpij += CalcPositionCorrectionCell(param_4, param_5, param_6); + } + } + return dpij; +} + +void _main(uvec3 dtid) +{ + int id = int(dtid.x); + if (particles_1._data[id].IsFix != 0u) + { + return; + } + int param = id; + particles_1._data[id].Next += CalcPositionCorrection(param); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp new file mode 100644 index 0000000..e87384b --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp @@ -0,0 +1,46 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(set = 0, binding = 1, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 0, binding = 0, std140) uniform CB +{ + vec2 Force; + vec2 Gravity; + vec2 Dt; + ivec2 GridNum; + ivec2 GridSize; +} _41; + +void _main(uvec3 dtid) +{ + if (particles_1._data[dtid.x].IsFix != 0u) + { + return; + } + particles_1._data[dtid.x].Velocity += ((_41.Gravity + _41.Force) * _41.Dt.x); + vec2 pos = particles_1._data[dtid.x].Current + (particles_1._data[dtid.x].Velocity * _41.Dt.x); + pos.x = clamp(pos.x, 8.0, float((_41.GridNum.x * _41.GridSize.x) - 8)); + pos.y = clamp(pos.y, 8.0, float((_41.GridNum.y * _41.GridSize.y) - 8)); + particles_1._data[dtid.x].Next = pos; +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp new file mode 100644 index 0000000..840cfe4 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp @@ -0,0 +1,151 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(set = 0, binding = 0, std140) uniform CB +{ + ivec2 GridNum; + ivec2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +} _66; + +layout(set = 0, binding = 3, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +layout(set = 0, binding = 1, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 0, binding = 2, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +ivec2 GetGridPos(vec2 pos, vec2 gridSize) +{ + return ivec2(pos / gridSize); +} + +int GetHash(ivec2 gridPos, int num) +{ + return gridPos.x + (gridPos.y * num); +} + +vec2 CalcDensityCellPB(ivec2 gridPos, int i, vec2 pos0) +{ + ivec2 param = gridPos; + int param_1 = _66.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = gridIndicesTable_1._data[gridHash].x; + float h2 = _66.EffectiveRadius * _66.EffectiveRadius; + float dens = 0.0; + if (startIndex != 2147483647) + { + int endIndex = gridIndicesTable_1._data[gridHash].y; + for (int j = startIndex; j < endIndex; j++) + { + vec2 pos1 = particles_1._data[gridTable_1._data[j].y].Next; + float r2 = dot(pos0 - pos1, pos0 - pos1); + if (r2 <= h2) + { + float q = h2 - r2; + dens += (((_66.Wpoly6 * q) * q) * q); + } + } + } + return vec2(dens); +} + +float CalcScalingFactorCell(ivec2 gridPos, int i, vec2 pos0) +{ + ivec2 param = gridPos; + int param_1 = _66.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = gridIndicesTable_1._data[gridHash].x; + float h = _66.EffectiveRadius; + float r0 = _66.Density; + float sd = 0.0; + vec2 sd2 = vec2(0.0); + if (startIndex != 2147483647) + { + uint endIndex = uint(gridIndicesTable_1._data[gridHash].y); + uint _195 = uint(startIndex); + for (uint j = _195; j < endIndex; j++) + { + vec2 pos1 = particles_1._data[gridTable_1._data[j].y].Next; + vec2 rij = pos0 - pos1; + float r = length(rij); + if ((r <= h) && (r > 0.0)) + { + float q = h - r; + vec2 dp = ((rij * ((_66.GWspiky * q) * q)) / vec2(r)) / vec2(r0); + sd2 += dp; + sd += dot(dp, dp); + } + } + } + return sd + dot(sd2, sd2); +} + +void CalcScalingFactor(int id) +{ + vec2 pos = particles_1._data[id].Next; + float h = _66.EffectiveRadius; + float r0 = _66.Density; + vec2 param = pos - vec2(h, h); + vec2 param_1 = vec2(_66.GridSize); + ivec2 grid_pos0 = GetGridPos(param, param_1); + vec2 param_2 = pos + vec2(h, h); + vec2 param_3 = vec2(_66.GridSize); + ivec2 grid_pos1 = GetGridPos(param_2, param_3); + float dens = 0.0; + float sd = 0.0; + for (int y = grid_pos0.y; y <= grid_pos1.y; y++) + { + for (int x = grid_pos0.x; x <= grid_pos1.x; x++) + { + ivec2 n_grid_pos = ivec2(x, y); + ivec2 param_4 = n_grid_pos; + int param_5 = id; + vec2 param_6 = pos; + dens += CalcDensityCellPB(param_4, param_5, param_6).x; + ivec2 param_7 = n_grid_pos; + int param_8 = id; + vec2 param_9 = pos; + sd += CalcScalingFactorCell(param_7, param_8, param_9); + } + } + float C = (dens / r0) - 1.0; + particles_1._data[id].Pscl = (-C) / (sd + _66.Eps); +} + +void _main(uvec3 dtid) +{ + int id = int(dtid.x); + int param = id; + CalcScalingFactor(param); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp new file mode 100644 index 0000000..b6892fb --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp @@ -0,0 +1,20 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(set = 0, binding = 1, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +void _main(uvec3 dtid) +{ + gridIndicesTable_1._data[dtid.x] = ivec2(2147483647); +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp new file mode 100644 index 0000000..abb90bb --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp @@ -0,0 +1,35 @@ +#version 430 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct Particle +{ + vec2 Current; + vec2 Next; + vec2 Velocity; + float Pscl; + uint IsFix; +}; + +layout(set = 0, binding = 1, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 0, binding = 0, std140) uniform CB +{ + float Dt; +} _40; + +void _main(uvec3 dtid) +{ + particles_1._data[dtid.x].Velocity = (particles_1._data[dtid.x].Next - particles_1._data[dtid.x].Current) / vec2(_40.Dt); + particles_1._data[dtid.x].Current = particles_1._data[dtid.x].Next; +} + +void main() +{ + uvec3 dtid = gl_GlobalInvocationID; + uvec3 param = dtid; + _main(param); +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/render.frag b/examples/Fluid2D/Shaders/GLSL_VULKAN/render.frag new file mode 100644 index 0000000..aa0abac --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/render.frag @@ -0,0 +1,34 @@ +#version 430 + +struct PS_INPUT +{ + vec4 Position; + vec4 Color; + vec2 UV1; + vec2 UV2; +}; + +layout(location = 0) in vec4 input_Color; +layout(location = 1) in vec2 input_UV1; +layout(location = 2) in vec2 input_UV2; +layout(location = 0) out vec4 _entryPointOutput; + +vec4 _main(PS_INPUT _input) +{ + float r = length((_input.UV1 - vec2(0.5)) * 2.0); + float a = (r > 1.0) ? 0.0 : ((((0.0 * pow(r, 6.0)) + (1.0 * pow(r, 4.0))) - (2.0 * pow(r, 2.0))) + 1.0); + vec4 c = vec4(_input.Color.xyz, _input.Color.w * a); + return c; +} + +void main() +{ + PS_INPUT _input; + _input.Position = gl_FragCoord; + _input.Color = input_Color; + _input.UV1 = input_UV1; + _input.UV2 = input_UV2; + PS_INPUT param = _input; + _entryPointOutput = _main(param); +} + diff --git a/examples/Fluid2D/Shaders/Metal/bitonicSort.comp b/examples/Fluid2D/Shaders/Metal/bitonicSort.comp new file mode 100644 index 0000000..246fef7 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/bitonicSort.comp @@ -0,0 +1,47 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct CB +{ + float Inc; + float Dir; +}; + +struct gridTable +{ + int2 _data[1]; +}; + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, constant CB& _19, device gridTable& gridTable_1) +{ + int inc = int(_19.Inc); + int dir = int(_19.Dir); + int t = int(dtid.x); + int low = t & (inc - 1); + int i = (t << 1) - low; + bool reverse = (dir & i) == 0; + int2 x0 = gridTable_1._data[i]; + int2 x1 = gridTable_1._data[inc + i]; + int2 auxa = x0; + int2 auxb = x1; + if ((int(x0.x < x1.x) ^ int(reverse)) != int(0u)) + { + x0 = auxb; + x1 = auxa; + } + gridTable_1._data[i] = x0; + gridTable_1._data[inc + i] = x1; +} + +kernel void main0(constant CB& _19 [[buffer(0)]], device gridTable& gridTable_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, _19, gridTable_1); +} + diff --git a/examples/Fluid2D/Shaders/Metal/buildGrid.comp b/examples/Fluid2D/Shaders/Metal/buildGrid.comp new file mode 100644 index 0000000..459799f --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/buildGrid.comp @@ -0,0 +1,71 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct CB +{ + int2 GridNum; + int2 GridSize; + int ParticlesCount; +}; + +struct gridTable +{ + int2 _data[1]; +}; + +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + uint IsFix; +}; + +struct particles +{ + Particle _data[1]; +}; + +static inline __attribute__((always_inline)) +int2 GetGridPos(thread const float2& pos, thread const float2& gridSize) +{ + return int2(pos / gridSize); +} + +static inline __attribute__((always_inline)) +int GetHash(thread const int2& gridPos, thread const int& num) +{ + return gridPos.x + (gridPos.y * num); +} + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, constant CB& _52, device gridTable& gridTable_1, device particles& particles_1) +{ + if (dtid.x < uint(_52.ParticlesCount)) + { + float2 param = particles_1._data[dtid.x].Next; + float2 param_1 = float2(_52.GridSize); + int2 param_2 = GetGridPos(param, param_1); + int param_3 = _52.GridNum.x; + ((device int*)&gridTable_1._data[dtid.x])[0u] = GetHash(param_2, param_3); + ((device int*)&gridTable_1._data[dtid.x])[1u] = int(dtid.x); + } + else + { + ((device int*)&gridTable_1._data[dtid.x])[0u] = 2147483647; + ((device int*)&gridTable_1._data[dtid.x])[1u] = int(dtid.x); + } +} + +kernel void main0(constant CB& _52 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device gridTable& gridTable_1 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, _52, gridTable_1, particles_1); +} + diff --git a/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp b/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp new file mode 100644 index 0000000..479e1c1 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp @@ -0,0 +1,58 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +// Implementation of the GLSL mod() function, which is slightly different than Metal fmod() +template +inline Tx mod(Tx x, Ty y) +{ + return x - y * floor(x / y); +} + +struct CB +{ + int ParticlesCount; +}; + +struct gridTable +{ + int2 _data[1]; +}; + +static inline __attribute__((always_inline)) +int getValidId(thread const int& id, constant CB& _22) +{ + return int(mod(float(id + _22.ParticlesCount), float(_22.ParticlesCount))); +} + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, constant CB& _22, device gridTable& gridTable_1, device gridTable& gridIndicesTable) +{ + int id = int(dtid.x); + int param = id - 1; + int idPrev = getValidId(param, _22); + int param_1 = id + 1; + int idNext = getValidId(param_1, _22); + int cell = ((device int*)&gridTable_1._data[id])[0u]; + int cellPrev = ((device int*)&gridTable_1._data[idPrev])[0u]; + int cellNext = ((device int*)&gridTable_1._data[idNext])[0u]; + if (cell != cellPrev) + { + ((device int*)&gridIndicesTable._data[cell])[0u] = id; + } + if (cell != cellNext) + { + ((device int*)&gridIndicesTable._data[cell])[1u] = id + 1; + } +} + +kernel void main0(constant CB& _22 [[buffer(0)]], device gridTable& gridTable_1 [[buffer(1)]], device gridTable& gridIndicesTable [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, _22, gridTable_1, gridIndicesTable); +} + diff --git a/examples/Fluid2D/Shaders/Metal/buildVBID.comp b/examples/Fluid2D/Shaders/Metal/buildVBID.comp new file mode 100644 index 0000000..1fde583 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/buildVBID.comp @@ -0,0 +1,93 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct CB +{ + float ParticleRadius; + float4 Color; + float4 FixedColor; +}; + +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + uint IsFix; +}; + +struct particles +{ + Particle _data[1]; +}; + +struct Vertex +{ + packed_float3 Position; + int Color; + float2 UV1; + float2 UV2; +}; + +struct vertex_ +{ + Vertex _data[1]; +}; + +struct index +{ + int _data[1]; +}; + +static inline __attribute__((always_inline)) +int DecodeFloatRGBA(thread const float4& rgba) +{ + int res = 0; + res += int(rgba.w * 255.0); + res = res << 8; + res += int(rgba.z * 255.0); + res = res << 8; + res += int(rgba.y * 255.0); + res = res << 8; + res += int(rgba.x * 255.0); + return res; +} + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, constant CB& _67, device particles& particles_1, device vertex_& vertex_1, device index& index_1) +{ + float4 param = mix(_67.Color, _67.FixedColor, float4(float(particles_1._data[dtid.x].IsFix != 0u))); + int c = DecodeFloatRGBA(param); + for (int i = 0; i < 4; i++) + { + vertex_1._data[(dtid.x * 4u) + uint(i)].Color = c; + } + float3 pos = float3(particles_1._data[dtid.x].Current, 0.5); + vertex_1._data[dtid.x * 4u].Position = pos + ((float3(-1.0, -1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 1u].Position = pos + ((float3(1.0, -1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 2u].Position = pos + ((float3(1.0, 1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[(dtid.x * 4u) + 3u].Position = pos + ((float3(-1.0, 1.0, 0.0) * _67.ParticleRadius) * 3.0); + vertex_1._data[dtid.x * 4u].UV1 = float2(0.0); + vertex_1._data[(dtid.x * 4u) + 1u].UV1 = float2(1.0, 0.0); + vertex_1._data[(dtid.x * 4u) + 2u].UV1 = float2(1.0); + vertex_1._data[(dtid.x * 4u) + 3u].UV1 = float2(0.0, 1.0); + index_1._data[dtid.x * 6u] = int(dtid.x * 4u); + index_1._data[(dtid.x * 6u) + 1u] = int((dtid.x * 4u) + 1u); + index_1._data[(dtid.x * 6u) + 2u] = int((dtid.x * 4u) + 2u); + index_1._data[(dtid.x * 6u) + 3u] = int(dtid.x * 4u); + index_1._data[(dtid.x * 6u) + 4u] = int((dtid.x * 4u) + 2u); + index_1._data[(dtid.x * 6u) + 5u] = int((dtid.x * 4u) + 3u); +} + +kernel void main0(constant CB& _67 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device vertex_& vertex_1 [[buffer(2)]], device index& index_1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, _67, particles_1, vertex_1, index_1); +} + diff --git a/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp b/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp new file mode 100644 index 0000000..1dea653 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp @@ -0,0 +1,136 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct CB +{ + int2 GridNum; + int2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +}; + +struct gridIndicesTable +{ + int2 _data[1]; +}; + +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + uint IsFix; +}; + +struct particles +{ + Particle _data[1]; +}; + +static inline __attribute__((always_inline)) +int2 GetGridPos(thread const float2& pos, thread const float2& gridSize) +{ + return int2(pos / gridSize); +} + +static inline __attribute__((always_inline)) +int GetHash(thread const int2& gridPos, thread const int& num) +{ + return gridPos.x + (gridPos.y * num); +} + +static inline __attribute__((always_inline)) +float2 CalcPositionCorrectionCell(thread const int2& gridPos, thread const int& i, thread const float2& pos0, constant CB& _60, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + int2 param = gridPos; + int param_1 = _60.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = ((device int*)&gridIndicesTable_1._data[gridHash])[0u]; + float h = _60.EffectiveRadius; + float r0 = _60.Density; + float2 dp = float2(0.0); + float dt = _60.Dt; + float si = particles_1._data[i].Pscl; + if (startIndex != 2147483647) + { + uint endIndex = uint(((device int*)&gridIndicesTable_1._data[gridHash])[1u]); + uint _117 = uint(startIndex); + for (uint j = _117; j < endIndex; j++) + { + if (((device int*)&gridTable._data[j])[1u] == i) + { + continue; + } + float2 pos1 = particles_1._data[((device int*)&gridTable._data[j])[1u]].Next; + float2 rij = pos0 - pos1; + float r = length(rij); + if ((r <= h) && (r > 0.0)) + { + float scorr = 0.0; + float q = (h * h) - (r * r); + float q2 = (h * h) - ((0.039999999105930328369140625 * h) * h); + float ww = (((_60.Wpoly6 * q) * q) * q) / (((_60.Wpoly6 * q2) * q2) * q2); + scorr = (((-0.100000001490116119384765625) * pow(ww, 4.0)) * dt) * dt; + float q_1 = h - r; + float sj = particles_1._data[((device int*)&gridTable._data[j])[1u]].Pscl; + dp += ((((rij * ((_60.GWspiky * q_1) * q_1)) / float2(r)) * ((si + sj) + scorr)) / float2(r0)); + } + } + } + return dp; +} + +static inline __attribute__((always_inline)) +float2 CalcPositionCorrection(thread const int& id, constant CB& _60, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + float2 pos = particles_1._data[id].Next; + float h = _60.EffectiveRadius; + float2 param = pos - float2(h); + float2 param_1 = float2(_60.GridSize); + int2 grid_pos0 = GetGridPos(param, param_1); + float2 param_2 = pos + float2(h); + float2 param_3 = float2(_60.GridSize); + int2 grid_pos1 = GetGridPos(param_2, param_3); + float2 dpij = float2(0.0); + for (int y = grid_pos0.y; y <= grid_pos1.y; y++) + { + for (int x = grid_pos0.x; x <= grid_pos1.x; x++) + { + int2 n_grid_pos = int2(x, y); + int2 param_4 = n_grid_pos; + int param_5 = id; + float2 param_6 = pos; + dpij += CalcPositionCorrectionCell(param_4, param_5, param_6, _60, gridIndicesTable_1, particles_1, gridTable); + } + } + return dpij; +} + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, constant CB& _60, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + int id = int(dtid.x); + if (particles_1._data[id].IsFix != 0u) + { + return; + } + int param = id; + particles_1._data[id].Next += CalcPositionCorrection(param, _60, gridIndicesTable_1, particles_1, gridTable); +} + +kernel void main0(constant CB& _60 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device gridIndicesTable& gridTable [[buffer(2)]], device gridIndicesTable& gridIndicesTable_1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, _60, gridIndicesTable_1, particles_1, gridTable); +} + diff --git a/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp b/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp new file mode 100644 index 0000000..dbda0c9 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp @@ -0,0 +1,51 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + uint IsFix; +}; + +struct particles +{ + Particle _data[1]; +}; + +struct CB +{ + float2 Force; + float2 Gravity; + float2 Dt; + int2 GridNum; + int2 GridSize; +}; + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, device particles& particles_1, constant CB& _41) +{ + if (particles_1._data[dtid.x].IsFix != 0u) + { + return; + } + particles_1._data[dtid.x].Velocity += ((_41.Gravity + _41.Force) * _41.Dt.x); + float2 pos = particles_1._data[dtid.x].Current + (particles_1._data[dtid.x].Velocity * _41.Dt.x); + pos.x = fast::clamp(pos.x, 8.0, float((_41.GridNum.x * _41.GridSize.x) - 8)); + pos.y = fast::clamp(pos.y, 8.0, float((_41.GridNum.y * _41.GridSize.y) - 8)); + particles_1._data[dtid.x].Next = pos; +} + +kernel void main0(constant CB& _41 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, particles_1, _41); +} + diff --git a/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp b/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp new file mode 100644 index 0000000..fe82782 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp @@ -0,0 +1,156 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct CB +{ + int2 GridNum; + int2 GridSize; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +}; + +struct gridIndicesTable +{ + int2 _data[1]; +}; + +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + uint IsFix; +}; + +struct particles +{ + Particle _data[1]; +}; + +static inline __attribute__((always_inline)) +int2 GetGridPos(thread const float2& pos, thread const float2& gridSize) +{ + return int2(pos / gridSize); +} + +static inline __attribute__((always_inline)) +int GetHash(thread const int2& gridPos, thread const int& num) +{ + return gridPos.x + (gridPos.y * num); +} + +static inline __attribute__((always_inline)) +float2 CalcDensityCellPB(thread const int2& gridPos, thread const int& i, thread const float2& pos0, constant CB& _66, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + int2 param = gridPos; + int param_1 = _66.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = ((device int*)&gridIndicesTable_1._data[gridHash])[0u]; + float h2 = _66.EffectiveRadius * _66.EffectiveRadius; + float dens = 0.0; + if (startIndex != 2147483647) + { + int endIndex = ((device int*)&gridIndicesTable_1._data[gridHash])[1u]; + for (int j = startIndex; j < endIndex; j++) + { + float2 pos1 = particles_1._data[((device int*)&gridTable._data[j])[1u]].Next; + float r2 = dot(pos0 - pos1, pos0 - pos1); + if (r2 <= h2) + { + float q = h2 - r2; + dens += (((_66.Wpoly6 * q) * q) * q); + } + } + } + return float2(dens); +} + +static inline __attribute__((always_inline)) +float CalcScalingFactorCell(thread const int2& gridPos, thread const int& i, thread const float2& pos0, constant CB& _66, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + int2 param = gridPos; + int param_1 = _66.GridNum.x; + int gridHash = GetHash(param, param_1); + int startIndex = ((device int*)&gridIndicesTable_1._data[gridHash])[0u]; + float h = _66.EffectiveRadius; + float r0 = _66.Density; + float sd = 0.0; + float2 sd2 = float2(0.0); + if (startIndex != 2147483647) + { + uint endIndex = uint(((device int*)&gridIndicesTable_1._data[gridHash])[1u]); + uint _195 = uint(startIndex); + for (uint j = _195; j < endIndex; j++) + { + float2 pos1 = particles_1._data[((device int*)&gridTable._data[j])[1u]].Next; + float2 rij = pos0 - pos1; + float r = length(rij); + if ((r <= h) && (r > 0.0)) + { + float q = h - r; + float2 dp = ((rij * ((_66.GWspiky * q) * q)) / float2(r)) / float2(r0); + sd2 += dp; + sd += dot(dp, dp); + } + } + } + return sd + dot(sd2, sd2); +} + +static inline __attribute__((always_inline)) +void CalcScalingFactor(thread const int& id, constant CB& _66, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + float2 pos = particles_1._data[id].Next; + float h = _66.EffectiveRadius; + float r0 = _66.Density; + float2 param = pos - float2(h, h); + float2 param_1 = float2(_66.GridSize); + int2 grid_pos0 = GetGridPos(param, param_1); + float2 param_2 = pos + float2(h, h); + float2 param_3 = float2(_66.GridSize); + int2 grid_pos1 = GetGridPos(param_2, param_3); + float dens = 0.0; + float sd = 0.0; + for (int y = grid_pos0.y; y <= grid_pos1.y; y++) + { + for (int x = grid_pos0.x; x <= grid_pos1.x; x++) + { + int2 n_grid_pos = int2(x, y); + int2 param_4 = n_grid_pos; + int param_5 = id; + float2 param_6 = pos; + dens += CalcDensityCellPB(param_4, param_5, param_6, _66, gridIndicesTable_1, particles_1, gridTable).x; + int2 param_7 = n_grid_pos; + int param_8 = id; + float2 param_9 = pos; + sd += CalcScalingFactorCell(param_7, param_8, param_9, _66, gridIndicesTable_1, particles_1, gridTable); + } + } + float C = (dens / r0) - 1.0; + particles_1._data[id].Pscl = (-C) / (sd + _66.Eps); +} + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, constant CB& _66, device gridIndicesTable& gridIndicesTable_1, device particles& particles_1, device gridIndicesTable& gridTable) +{ + int id = int(dtid.x); + int param = id; + CalcScalingFactor(param, _66, gridIndicesTable_1, particles_1, gridTable); +} + +kernel void main0(constant CB& _66 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device gridIndicesTable& gridTable [[buffer(2)]], device gridIndicesTable& gridIndicesTable_1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, _66, gridIndicesTable_1, particles_1, gridTable); +} + diff --git a/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp b/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp new file mode 100644 index 0000000..1fa10db --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp @@ -0,0 +1,25 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct gridIndicesTable +{ + int2 _data[1]; +}; + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, device gridIndicesTable& gridIndicesTable_1) +{ + gridIndicesTable_1._data[dtid.x] = int2(2147483647); +} + +kernel void main0(device gridIndicesTable& gridIndicesTable_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, gridIndicesTable_1); +} + diff --git a/examples/Fluid2D/Shaders/Metal/integrate.comp b/examples/Fluid2D/Shaders/Metal/integrate.comp new file mode 100644 index 0000000..6cfd416 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/integrate.comp @@ -0,0 +1,40 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct Particle +{ + float2 Current; + float2 Next; + float2 Velocity; + float Pscl; + uint IsFix; +}; + +struct particles +{ + Particle _data[1]; +}; + +struct CB +{ + float Dt; +}; + +static inline __attribute__((always_inline)) +void _main(thread const uint3& dtid, device particles& particles_1, constant CB& _40) +{ + particles_1._data[dtid.x].Velocity = (particles_1._data[dtid.x].Next - particles_1._data[dtid.x].Current) / float2(_40.Dt); + particles_1._data[dtid.x].Current = particles_1._data[dtid.x].Next; +} + +kernel void main0(constant CB& _40 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + uint3 dtid = gl_GlobalInvocationID; + uint3 param = dtid; + _main(param, particles_1, _40); +} + diff --git a/examples/Fluid2D/Shaders/Metal/render.frag b/examples/Fluid2D/Shaders/Metal/render.frag new file mode 100644 index 0000000..abc4f9d --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/render.frag @@ -0,0 +1,49 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct PS_INPUT +{ + float4 Position; + float4 Color; + float2 UV1; + float2 UV2; +}; + +struct main0_out +{ + float4 _entryPointOutput [[color(0)]]; +}; + +struct main0_in +{ + float4 input_Color [[user(locn0)]]; + float2 input_UV1 [[user(locn1)]]; + float2 input_UV2 [[user(locn2)]]; +}; + +static inline __attribute__((always_inline)) +float4 _main(thread const PS_INPUT& _input) +{ + float r = length((_input.UV1 - float2(0.5)) * 2.0); + float a = (r > 1.0) ? 0.0 : ((((0.0 * pow(r, 6.0)) + (1.0 * pow(r, 4.0))) - (2.0 * pow(r, 2.0))) + 1.0); + float4 c = float4(_input.Color.xyz, _input.Color.w * a); + return c; +} + +fragment main0_out main0(main0_in in [[stage_in]], float4 gl_FragCoord [[position]]) +{ + main0_out out = {}; + PS_INPUT _input; + _input.Position = gl_FragCoord; + _input.Color = in.input_Color; + _input.UV1 = in.input_UV1; + _input.UV2 = in.input_UV2; + PS_INPUT param = _input; + out._entryPointOutput = _main(param); + return out; +} + diff --git a/examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv b/examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv new file mode 100644 index 0000000000000000000000000000000000000000..59433053d9a5ef618e42d12a808937973cf618b6 GIT binary patch literal 2700 zcmZ9Mdv8=l7{-V0wpLICp&%-COHo0vQn`!37Ao6ny@L14_Ox5;>e8BJsr|!0jNe90 zO#BW;Ka!Wk#NY4iyy2viyqV{D-pkCqGw0MsrXNaaSK6J%(>G~2_NNh)l=h^ta&9f% zUYfZ#=**luf5wQtsa6IW^B6Is%ysN;(cRFog6%_B+m9I+RwH)={wu;L5!Jj=ls)*` z7CNzY|Lm1v%x<*Kpxa?0qQ;sjbvIT^Z2g-zQn{Mgv53#rW9|8Fzl=SAuRVLY#7DrL ztWSamn6-~H_cp&P@kwyE)a%KE_t&=8`uEmK-$nQ~qI-nlH$Fds3|@x1M{@@8NwDPiOvh=0}(huq*d9MK$Y*C&+RA2Ie>n=Yf}+ zF%7tTQ`4s44>OyqZvLap4a_|lzn9rq^+T+*uhdhT#Ak2hL#>`32U|=13VyjgQ#+Mw z9|SwBJ;AKc{O~^sc5S(5c?NSB-(ZgO93*ol>n~+p?wQXsyI*UYcY)bF^`o4@dlFBA zy$`u@uP~dV=Kkg0Nz}QTvAM<<%)2ma8!unUTI}C%$e7cd#yIa#E$jmC#QMHBhxhb1 zivJaBdN0dZtReSpSGjLnewBFVRm^={quO%j;qx0+(`Oxd_}0N{`m8DUkKz8l!hDPB z|8b&ASnOTh;aj>$u6y*q;q2Vcm1jx$bx4`IGn->5zDT^`2NH{iahC$eUbJdZr`p8}ij8;*RxUG>O63RW}U z^UKX|Fi&Cj%{OoU81peq{rkNBbk@wbUvgvJ+q0Nj@aMqp@ki?Ge;!kFe}87}1lT-% z*0kqpQR4-$HGj|Xr@(4)M(du&%!_qi1RJMrP47)DYR-Z^d+?XR#>Ae^f{juCnRWfH z&tdA`jWxY{wW#qb*cx+m&F}OY=C^CSKI8muuVebG;d_xsjXALW3Vs7@%^&jFZ-R|c zk9Fq3v5x*&=QXfC*J(1#!`A}4wm#paJba7bCZ^B#DffN&{k)Eu7w^N~zmfHbS%NpQ zYRsE(_1|(Y-vXPb{uBMO&fA!}_iE4Fk9~g!v)2*#Zq_31kKCL0(5<2Wm%Pqxu)6jB z;ymv&JN~7%zB`zj{|RyQ-~+H{k2{iw??bRXi@wVLAwTLk8|#@q!o0u8{TOWj;_f~H t$K4rgO*P}9#;4%woA?Y~eG{Lq+b&Pi%HRIx6`ahxCr{Dkp literal 0 HcmV?d00001 diff --git a/examples/Fluid2D/Shaders/SPIRV/buildGrid.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildGrid.comp.spv new file mode 100644 index 0000000000000000000000000000000000000000..91ee3dfcb11475023221a29fd8d79ac67bfa4384 GIT binary patch literal 3532 zcmZveZC6xf6owBDAYr8@YH1Q=C8f(0d?!sDb0CQ@jNYtigb}9(#+ZSWDE-ifeod=? z*UMT(*EQ#CJZqMRwb$PF{p`2rdGG5JxXglh&8=oF@H?mx*t+?!g=7G{vbK!fle~&uw(*hmwXyNEr^9fIt=b~0?LUexUQ60L)%LN$8FKXU zC+Gg;qMu&&BJW;jE32*ALcQ87Hr7^JY0fk7<+G#Y?@q2@Kd7u$meKa~el7;HU02Sh zxVE}lC5fjlXWg0VL$d1E+vcn}U#&M5YOO~ppZa=hv(1Hip7Z>tnm20RnYAa-^pf_A zZ5=tE@?xb`LF-pUZ~K)`!*nM%Q?9a{;n{7R?(tX<+B0^66IvLF0hn@{`K zzaK1@=Iv?MmNQNOWWw!hPUYF2= zRB^Vf72IwQ;2=VnfdOX&Uk&!VV_6*Huyb=GrExR zo9x4wi-_DG#9Ts;Bj%C+ll@;t`}~8yYwsc7Bc6wL^z#9BXFoUK_(N%R!r_u@L@ zS_!c)&+$XVb!XnQ6`Ok+{So3m!oMGTCgU4f{Kpw*ymujv+@D~Z)3sY@ajf0O*55Vn zM0}BNgf%`ze5Z_YO&{-7egYrQ_%p=b+*90zv-dgTnlaw3IL_*8?9Q|L22S3-to-#tJ~(}a_cy|u9AqW&6o)Hh%3OIA5Qi^q`XVS;LGTSsVSs9(+7=3kAf;%C7OaB!ZhG5<8V#w_hK?Wn;P~CY z#2NoRUuJZCo_p`gnVimccYnL*w|n-SB%>pzCbDc_wm+N9dRe)qvJsFhJCKbRy0Ln1 zb#8aip1W}IO*0N=qeY@QhsYU2*N{hfw{M`0JPR&1g_svI%H9$D-vP$RNcWAgXFsv6 z&fr1b>$aQineO~m^G36$z`GES@h4}?Vl#V>7mIj}dej}5QRW{(uiVUFoXE6YYx&Ng zyV>jPu53T<56Zlw;H`y)qQ+_LvPLbN#r`4R$seKZ^Nk|5w$s@weFu@o7sGGO7x_y? z{@u=xMV(r9DdK4P-a>D7dd0{+(>469oo@SnzR~MoJ5@1PtDO&Ww0ZBK|IaIGjb$rf za+Tecd5sDyI}n@5KCbpQ`|a*#XE#a}wT!u*)SAfd<6A4urzBjpY=hX7c~2I1zSV0r zdfOX$ui4+*-pmKx?SAtXdFDSvr}O$Y@o$q6cXSHe)!@#>d@Bw1drv13SNPMVTljBd zk0XC4BnEaD%bQ@)r_$|n|l&MM{8moNJFVIDwUL9G9M<-dqN zOqO}mu*s+WJUP~{Blc9s=Yv{^_%yy})fPAZX>=WN9m9VHEheY?52DQ#cmE-@wZu=5 z<-E(9HTcSy*i~}%y#w=S@y$Dixa58TZA|ynF`2e zJawdB^>-3m{9pRkjx&y6dv=$f3D4-g_?zMxpF?{$p8GVS4u5vY_je+V`#X{Pi$ng> zknitAn(yyK>aPs>jmmdcXR&7x`#ys>b3bFRA>uz(b9lYNdVv&CtTTCwJq#w1e zU|UPvxN{YYy*IF}waheqo6Csb6~4_INX*N=Zz9I!asK+ja|_#Z1kbzJ!DB4tc@N*X zJl>mr@Z83>$KYvV2amCo=Y4$R@_66+!SeyOJ%+!EZLQz9gFEQEh&keauzwBh`kS|4 z>_bG%cdi{gpJ1oF5Af6ZeTwgQrSmh_e(g`(+`r%pUUl$(hMn?$j_;YG-!HJu5x3v3 y&@THm_9Y@_zuLk3HFnDT4Zi)x+iGEp`L<%_Id+`E{p*Nl5I5g5L@jlEFa8HvOU~H< literal 0 HcmV?d00001 diff --git a/examples/Fluid2D/Shaders/SPIRV/buildVBID.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildVBID.comp.spv new file mode 100644 index 0000000000000000000000000000000000000000..520b65e750dd79934f2579ebabc9fade08861179 GIT binary patch literal 6956 zcmZ9Q37l4C6^Ad(42vN9rX~$SMTsPeppiMlAOlG>qLA4#e9U~(m*afX%m+bPLSeg9 zmaSH{YnEl%UfM#NZMImp@B6-)T3XL@?>+E(ub$uepZ_`US>F5J_nRMW;}=iNvT@l_ z*`(~aaasP%%*MlH*_3QTu6s9~w`uuUYjF9>(@r*ETGp1w88e-j4s<(mNv$!gV>XT% z@NzQ|<3ie4J0AN6hz=ssbsh9D72mr0Ky$FZeyCY%ZRuUxy>$0QE7w5h@l(*FJ9pFy zyCa(o=(q6F-MdaD%AB*{23w86BBm{y2N9jynpyK% z7$dJE=hHt}Yt_)k^`Q^P^{;?w%lhF4a({caz3?C2Gx~((N>+J_MA5a*Vi0t zv>MI)Jh*@I_us$voO8ad+g7l$u$^n`wo_Qqp1F8;p1VKqJx>Og_C9hXW(6?s{n_Y| zX7Rqi1#S!v+FqVR+viA*_Uv}-JSLyLJ)23RN6ycr{5hB$D?x+oI9Nyg=s5#ok zb*!_`xx}-7<^Hs1^Amdm5#1K#SzM?rIKpyn20A$R!?;eQw>+P6z0~PJpMdy;y{Itr z^_xnZ){%=klfc%Ii#kzTt_^YR0wm%*Q*sYLtoa!91f&x&CeBmTv}ZYU8SmLjd)BsR zdvg5{+9%cr7d=nK7hJVx*ILK5_G~@B>$JZ(Gm+o^sihjr+*r{6Y(+MbhMoH%r^sVE_v5`f7^$={<&!V);SSdJs{)hMPpo&p>@U_Y0b4Wdlfcds_A0Q?Lf%I=+ML!<>~SsH_4?M7c2CJg zO}|}p8h>`_bA9+Xl$?8X4%iyzTu2qa6}b?7&V4*n`+ewN2;Pj=w#H-8u9veX?Y@%J ze=gYi=7}E82Yc@HZ$-Oa{~XW-XydGXA>tGE6LM?cS@wQ0*cji(LXh8#obN6Aw{O22 zefB+!*0ydFZCyG0*7h5bv!4BoAN>*q%KL>|E9~-+r`yYe&Baz&vM#{gT|= zcgr&Rb=?(+oZ_tJIfyoov&FhA!CBVTrT;3Be%D8ytHJI~j2ZN3`w#+i3fjbKkYSOXOnT?*PaA zzK6-<{5>kx&jRy*{vEqs+q*~3`KF`2drm@!tYWqKeV0C-nY&=jX%6oIe;a+DVf)*c z+Wz)wpGmwm=3ra*QDk$icKGIj<@A|HJA4bka{A1veMZ^eG2q$cmcN&YPDkwBeVmH; zc<-%dt@p|l#5s;e+qZX^e-qlL`NiI0-Q@7TIu>z__n7^-$HyV^mlLCXe93v24HHwm z!-98>_5%@f?<;%jEV+f~s4Exoi@?SoD&rS}K%;Iz+W*m7x~a?$5X`iQI%cXtFMW3s|_PGj)K2HbBN1tba<ZZ&PK0CdXRO9ed_B)CFlO!LyUGEF_&w_vF}d#(Rd`4|!vq*V*k^9eJJ!wszc? zmxASDebLwb_tV*9=F7m6{QBIVtI)4P;@#Pgb(}}u_?YL_;OacDfs>DUUJI6U zo;W|R1Ix#6f&IQ7NoRZmc6G)#!pX-xZvtP9$Om@~SU#Qc&DiF@28kKp0*)E&CuWp4 zK4yF?xH{w8;N)Y*w}a)<8Q*~|zXb7YIpaH#bjItjt24d}PCn+j9(*k#AKbga^68B4 z!8X6XnDGX1%xFI`qrCAk<9or?8Q%vdA2Z$vmP==RKeqe=Bxd{oIGyo>*wq<71ScQ! zd>DKaA|KpG!1C#gAH_DmzL@c2;F!^VVn%u6W5$nzt22HAPCjP*Bv{TF&F@`t3ljPD zZ$SJOwXB9vsZubCZm6Z$aizM??wL@ zxgGIrS;zM;XI#|!3AkG8r*QI7>t|p&dyKpA=V1Bb{TK880=t^)mvHit>sMg8bdP_H zEpM)vl1F zbKM0eAGz)Z%Q;K*eGgbZ`u-0%`nHaJ%NZB7{tK?wx))A9YTXBxOZ)yGw!HNtmw)6| zbB%+Ok6h!ya;}TB*@i8jzJqpb`Iw~x9J5&4S>%k1`V+v_`V-;gqy8kYTsq59*z*3H uM&z0TuI8EwCm*?v2FsMX+9>= zgd`?0O^hanMAIzkJsMLrF-_BZk0p;v-goYtWw?(6=lQ?yJLf2v4vy=zTTb+xRe|S z?j&ekc}E8jr7k&~mcGkdDwj{$R-Qh$++e?Aa%dJyLU#zco}n?(X`Z;u2Cr&mW+qLeC{chmew-P{_WyY zG8x?LB2L-1bw=thCdZ+*_Zv|iV-9U^OJB?8j1yMn%?)Y3GvJMjlXO?aUM}V;=3B=< zHTOa-)~=3qS9f2QZrtefw{@&-X)kZB4CHagmvnTqNZZkhX}o*4wAYv__Z%$g-#2z0 ze0^_E*Y?@_G!9$Rw7$|?zG8b??|gXk>}l!z8^P|W8wf-zCQE4Z^If(dFC{DBmHw8# z{uSNr9aq_BNOBgMd63F&9hGxiTD#IS;z`x#)!g3F-$HBr*|fFsEGND~vKF7kh}YBh zvAg08;P&1szEsS-)my<84ut#J23}2R<+gp}t9dnU<(zf*r|YZco~-J)iY%^GUE0_E zIKQK-r_DuA^=i%ncy*<%E3dg1D=N#%EI~T96K37gu`6Kbo-P7--`q<{XFe`HL#3n- zt)tuJb4ATpOs=GDZfH2Q$_vR3z^2q!N^Yce{bHdVOU)}k)=v09Tsl@Uxdn}ThSjdF z-!9s|@+LAAlH0|pUQBk=R@%5ZQ(j2!qTQ0=J>u!r{ttk!q*(huuAbc&Pl78Hto{O+ zF;ncqSNFG=6u|kjHZ*v)E(gM`zd5_fU6+F*hMe+|HT=-v&C^r=Q3HGiyp$XRkDM-o zImSjz#>dz2iSYJb7Af@?l1Xq6h;usy9A{uIxVt&BS932wOXt`;GaYNu0H2le#sNM% zG}^ZgKU-dAcz<+M_!9h1|1 z))_Y;r}?Z?I~HqVZ>-O8o0yBwy@{At@cJrE`I%sIM$PfuGM|_E)Cwssrq8f^KA&mk z)9+l`XuYC;Xtir&T^NEiD#e9yW{}42} zz6o&mMol{$J}lL%^&3ANu2#)kO2<~SpJoovpZ(S4qcU#4de-m|TCs8Z>Z#yJTKj}% zu6p{8&2;0er=E5K=4Cw-Y1N#cW4Vszsp`e_oLZ~;Nm>7?V6WglqXuE)sMzyt-J0w2 za^^Vm@R`^6xf%ETtC`pNEx^2tUr4J@eJs_UN$VKnz@B-zaVu$!Q!}sJbFZd<4XyLD zt{7(>IQ-{kIPbpz&Gj)}{bJ1PWcI>bmtbmI)c3dO=`|V-kE7JhA;cQd-q;*7?J+%XTL49juZL4Ut*oSFXYC(me%}g=9llxw8(!e zd=O^*ZJE#9u@3ixhd(iCUc^5Lugm!#1wZ=8BM(RXWAG%oc45Rn4KttlhcS@#SZnw% zLMvwe4Vk|_-+R(F1qf`DC-NJbYWgYWmD04_^hWrq7)6Iaz)E;K}4xe~}rThFQC7<8L7^ zpIN8U*ZX`#4j%}1P6yLEmbFg7j0t`M*tna>ssBXG81-)YRQI5&zeKK+(4z+LKYvqs z`5dwapE>@X5|5(wU1<$R&@N&OwNo%-e0Ered8T3N5yPKrX^i_7F*CsG!DoWae*<~V zF$*(B{n||9r$w4upY_Y5e&;&|^YWP)_4^Ks`hE6`U0Z*@IhSK-JrmA#4rU&?bMyCB zo;T*W@k?^t_+>SGMUJlludU(hbKIKzoe|?1?{AGfZ_V*7;JzBppTYz2`u)2-`kVij z8on#X9q;xUzPpCsRm1nxaQ_~U@yzGnIb2HSdSVn(J%G#cl=bznYx&wA(N> z>+?R5N8YQz))%qc!D_~O56O)+&nq!CeLFJ#Jat|TwqDz3hU*6Dh%G3u_3{8gB_%_WY$H-XI+=kaymotS&2FTXaghwHcRFxod@?(^N* zxNig-qwfCPOzUM0`rd@8S%cX5Iaj%Q)bwU>zNWXpqb7a%n%)Z6Z~eE?z731BrtenF zyhCW;j_G#~ZpnD~-T@zt1%D^lnCqFFecy!{qi+9QnHKTyhO3R>PV`=U4>kx}jOnwF z&qa0n81r6CE%+Vawb(=1dcF^A%pF+7ydP|z2eOzCfYmO>c4yiL!S>bXKHN#`+-ptvOQUQg!RFUzZF^|d;<@-x zswc^Vw2}8=uzK*%ft^#-`+2Z2>Jjq=a6aaXaO2b?=1bsw%$MQDsYg9u0nf+s^?Vhs z9(le7b{b|fk6_lR&-HsC(;~-r;A$hd z!;JqfHV7My>3a-|I==^g6f-9Je;<4)7BN2n8>b#IKLqDvegrp8J#tx(XIi^3}lfM9WV#eyb0E@bQ z3D)nJPtwZu4WWGs+l9qB_!W3Q_B0m#e+^a-{u{8h$J+fCY>axu{0^Lt`90h?^_$sa z$NmHM6s9lw{SoY%tvh?A`Wx67_2~0=@NO(( z{sA^lJ!kw2Gdx6+zVpPK!zPVwOMJudIfxNX$;JlGoiT_eu_t{D&a z8jJb+Wfbi=%o-=q`g>(4?F*Q5ldqxmrVOs%Z{B(CH}Bvh@%zm?xbc4P&hwEq-0$6a zzu&uqo6qmv!OiEl?mYKfcb-qoap&l_?(oa~-W}Zhe)GA)zTdmUZ-2jc z2RGjD-NB9bd$-&@^1go&bKT=_mY2Yt*aA$SdmsC`7p&j2;J(P?J|9Hh`S5j^@pZI= z(W-hO-8XsUDuT@su_d^B8}~pRSWSz+dxn7Rr_UPFy$4%^-(-fT`o5Znlk$K}@1Z8= zJp!&~f6t0MYC0HfPJNyedH6?#(3`3qyJ&x{Cp3G=jVF_ zTuqDn`AD$+^qJGMsTT7+I@R}|?`SmjnC}>{n*BY?@~CMX*qr)2$MWz^02{B*GcAw# z)`Oj|ee@aQeW5O%NP7u>>+}pB3wJ#lz~=Yd9S2u8-g`psb7V5D&j4G0gG`}40rS#- zV&=DvTql9e5#yZb^jKXACEbn{#I9q`#-0sU zvyZv-skuk)muqMnHLU`>hVk9*9I#rv8}+H>&wzU&_xzgAHu9|oy9beP4OlJm=~J_R z+!LP7n9o|UHT!NA$GdefbDNBL9gF#HokV*aW(~*F#<#%p=;vAlZ>`}S8TVXnA?|#* zdE#8wgV&>}``y9&`U0?;YifRZ%=tpFn(vwov=?KY^zmJz?;_0DdiW(^{qeo!Qm~q1 z`#bV7u$O0FUo)oW9L0W%7OTbhZD7Z@Hs>l2Uwh^Yz7cMWv95#MZ{Ds~CwM%qdb~F; k2gh8@m49!R;riq3Z33&AC;DF@C#P!>{T*A~`1rQ|e_Hlr)Bpeg literal 0 HcmV?d00001 diff --git a/examples/Fluid2D/Shaders/SPIRV/calcExternalForce.comp.spv b/examples/Fluid2D/Shaders/SPIRV/calcExternalForce.comp.spv new file mode 100644 index 0000000000000000000000000000000000000000..b2a9f27d7c09a2d600c74914a1a06117cf892c0b GIT binary patch literal 3048 zcmZveYjaZt5Qa}LDYXRzxzrmqwJO+R1t~WXFtj!lYOPkmdpIPqM$@FGX$SqL<7fYf zGyXjtXPohQlCy9odOAD1`|f?uo=iuEXAdOFP%@E>CqE?p=V&qvk|dMKSdW+2HrEz* zyS0VIi|6z>oQ(7w^?4ef0(=zOP8%%^4d@8C*ilHopb=t+(fS`O~Pe1G5t2f)# zM)$`)-j8>xjooT9kJ;bq?n>hk+}h5wlhxiuwzim0rIvP6xPHgrkNqm=K}M3(U`x09 zG0HwG?M~Gp%p+#M}?ewUdbeC)H_Y1G#Dw!y}FKK3cL+4!lq6YISWR$1)V zej9nGy-StRWDEJpKF53SYpYo)H`@=>=4xxdT}``ww;vzvely;F%)686&%uvz0q$*v zoYqx7K}_reHJL&`02kAzi0ehP(ZoJsFO!qxRzsHe_&=53_UHt^)qi+6_J*}T_3OQy9|L4Knv-Mtpo7=eK zNYn7}do{B;+iN}hli5zA$6Rvlk~z=7&8d$wsQa$OBKI7!Gw45`+03JUKGXZV9ZjyH zJ44vZnXcV@@;eYuDe~V%iFde$Y;AWH_K%Q#hw9FLABujy&17rZzq+@j{Vd+%-j*>z zm-rs6;a+{L{TE09ZVqd0L6KA4--Nuro$1%{7yk~j?hUdTs~$FgV8yhVM?GvUWHD{# zRKJ?_*G8U%#Q$NUD^T<CP1!R5fJAB?m7XO(%#+-wU6E88n_xl!PoZnFF>{%9;_k}6E<&*@dp1YZzLeRc_GNT2^F;g=*$e7 zTYe8W(6vXto5+JZe-}(V@|BPypSJvLw|aj1_IJ31{(n1M2Gvi)gu zc5hb9Ud!;?kXZiRtY*0N*O1?btbZ2@KYv5kA#LIJ0kVGHblz_RUE8nJH10#lSn*#n zY%|0CJ!*4<-pI$0_(8rv^-rK7$lI{TJCHsx%RS_n#aY#j3Hzs+{XX0~GUsQIc+BuQ za?Bcj_WT8;E&RUB{G2uK_Z7OfnDuL9W5r|E3bJ^9)(7a~F>8t(v)ZGxK7eADhse%q oZD&iizooX_}Z`Ok#{JQS<-z-OU_ce16-#*4q2*v(LHb-us4E(|^Q(EbEsI z$_8fl^~>_lkgPvSmTjEX=e}|7{JDGe^tSFbe!qQ8*et8b^UT?toH~3hw!BpCFmM}z z!RTs3F!RD{7~3EIAVM7(>A1RV0JXEMv8}hUyWBdjvuE_mCF3S*JbmSRQ^&3~z^ysIrPN;TSURiJ(%ac> zaTbE-f$k)Zk5H^Lo?Z)A^rQSnS6Oa|C-=&wpxOSGBkf^;|>s{@BbVOWIm` z%PZTCF140d^sKMr=_&Rbtv1j19L(%8r@}o`GkW{xu7@w^>TExA|2~a_n>8-z=_;Rk zX5Q}%c+>v-gX^+6VE5I-07R+H4#lVIox1_A%Z`Ni^p?7N=XA8To#B`P*->cLK`OVj z^&C@L+@7B|_r#c=rq)t#32**e+%~s+$3skD)pN^kLtL2uye$@O-U(wy&*3p~qx0=`Jh_*A>P4})oQ}g;9*Vf+I z!lAF|b-C`fnb*_e)a%E4ZRYgMD%+U7I)Hu5=XL?^MP3Blw-2-Y|croXLaBHP!;i1UETjP-pcwP1w ze8aO)lRb{+b@Ez1iSOo)7e8$-ca?Zuj>m!Pvgg4YUens_1+=EIllH^q`7h$*8ovbY zyT&!y%kb{(KMpRrD8fm#U5FeV) z7V`!tT1}yCk!U_w%-bf>eD165kZ3;p)rKdU&wjO$SR?1a{+zdwwYa98F+af@Dm3Tw zz}AeOk1Ar;^J6`=njF{W&-ApO&v@$@cg^wKjq#yX?@c^+&5ONJvyNKqgIes3``{Wp z%*L&-9)foj$hZv~3=Y|VJDr`*_ZxYtHa+Y-J}u2=due-K=)QrB-?S*|%wvxa-@ zcyr}L3huZD_V6Qkv3bTC=wNHSV?wi*`i^+}+fc8~ubVwM-abd+hhu*1b2MJfb=$*k zcZ{CKdR;1j_8uwkrZ8s7D2_VZ!pJkyC;*Zj!^_r6fGuIrwO`7wVM z-kAE1bUP349K*og3v%<0#+#>RUAgyzn(+mA*Jocb&mwTdk1ue<^(D+tJdBFo6V?rL+$hXmy*G^8tM`E1Ue3ZhubREey${qJ=RFW>^sLJrcOKpxHOI+4 z?`qNG8gM_%{0obiy+!;MaIDcieG2mojN#mw_Y77l%o@*u)m>B6c%#sak7P3Ew_oGd zUW1lptEWW%WpJMx=HCi7>^XA{a_>>KF^pEf7vk9p^9|@6*6?G$Z=?83wuWca_lkW6 z_q~yF-y8CS$ajw0;m&(7)f_92*qvZCW7d&J>~64{F>A_wgWBIc;4ygh_gK+H%--D_ z|9H1!yB0NDB$ zQr8-jF>};c7utbfYa6qFdGzmkx5xZ!jYa?46xd$3EwFpL9p+kgzlDm7a`UHF z@i_@M|Imb+Kevh>neYYRg;l&c;f_D0ink}+p8Z=Y`m^38ReWv1-BYjiA(-p;EV?%N zRPLc!_(p67=Cd@tcC(9^V_d&H>Kz8QjxpCRkL!0h*nDH|gM1Qmc`c8?JX?-2W{!KQ z9^;Pzt4(Ev_IE7i^V$B4r9B=8H*OvGSngRn9)AMnx)x#9_IYz6rXD#bfva;)MpM6t zI*w_=9H;I)UOP2$GrojH?~B3q@c?66@okt|3-%WIOYwf*r5|HUFg4GgIO>$a=^iYD z#~v8-c^A2-f{ojc=SCi}YLgI8f*S&AkwR5q36auCXXCB=I305(JRNft+#L0_#d*0JY%OEk#l!d&OLDR)FbC!a5Hv( z(f@s5wRa?i5+qqp8h_>W?qxp;?stk66^)-$G--XES_`BJ>~^wHnP z!Ohq_`Wt}%1g56>PLjtv$^CD@{OpLuJ84*fhv0WA@Id@<%=vc4$2;kMu=B(_>47T# zV8RE3A4<6OAF1MxCERz&<5m30D*klB?Qvbgp94Q%#b2o6FDBgaFIDlE6K+4RR`J&o z?)cXW?)4Z1{uJhZdM({6dA#F34fbr$#KvQBCO!i;<{0-=9`zmuTgRAtDUWfV1)Fco zeU-;M?osfTm}88Y95?{P*zotM?uG1N;v$KgNHAH?EIbKL%SP_)oz0_At+gR{T#fbJX1%`S&qv zTT2{ce+ITzoROb{S7DwJW9i=f0&YCsk-r3=jhSm~4QAZ|_+MehU9b0`JYuhahbR2k zU~}Rf`5W*nn7ZS=XVqf-Z^3Gta1uR>zr*@r(=cO>c@?kj7;}D)sRjQ7*mE22!#{$} zQIDKIfgSTiv9EsytCg@%7TRCHjy2|){2Shn`)=&7n40@8j@o|*JAd$hfSoJeqyGe( zqwacM!>gI^{{9Pd58|HtH+Tm21{O8n1gi&s3v8Wuuf7d7M?G@h0jD|d!p&2^x_Afu z2kiXDV%&eh_8;%5_rPje{C)U7*s;c9+y~$%u^8utRMTQyKQzZ0b1z;Og010R^ara2 zuYs?F`%X7r3s*DG@y665zYd&^uP=Cf*BSsfM?Lx;2!6VlqaMEzT+LkPFs9}`9M@kj z_x|=g=&d&jKZsoKL%%o3H-@XN0b9?QTKdfJewIIrx1K)w+k{+uh&65sUzhOB3Lb0R z9Bz)f`{tSkgRNuC?;7&>uF;?PE|{MYnBO&q;74NivMb*25Cif4-r{wTFTneoOUnJ- zCAj(i_LA}u2_H<{-(J$Vzq_Q|-(6Df?=C6#cbDMSo0xFt_cxb_%l*wIxcwho#r@4C z;_ksARovfRBJO+(tN4itcl^l-cmC#tJAQG(&2LS(`AZURzQ4W1{O0@HOK|i3?IpPR z{`L~we1CfhZoa?01b09D?M3cc+X%k}+;b4`i7nyro-pQ_j&r#c+_-1lb0Uvt+BRUH zX^V<{&*Zjf>gIcni% zXRvjQIfrLg&7MZ$cfsr_);tKjYoT{CulqI%uI6~px;%Q^9c)cwo^yG`#(>Q?=9!oK z{m?n~1aF3Sj4^Y(Z`5P_-r#io`@qxnk1cqte;nK#b)OOTzAxB1#+<`@O)b{HAGmt` z`=hDH`X_+Z9Phm;j~)*ITho~Lr95H>g3UMPJt~j&9|U&&jxlDA_qn>*d$>eg?ENdY zXV3doxO+7L>|CDdX>fJ(55dcQ_D;w9ywm&rZ6VjcYNgg;V8^NZ-9kPCb3W%6N4>+r_CKC+G3OC*b^G^uB)7jK@s8Kq|55m(F+axV z7jb>mItFZwnD1Dy{ipLC2Um~&7l0k7ZvXOmnDaTmIO;C~TR-^mVE4;&=-EC2Y>v9? zS%_CN-}p&jHDf2@ z=2Eb=jQL!aPhvLbI|aNw-g=HReiUo;SyToa_hT)!WpK~!`9;pDU^R2CZ%i$EUJj0Z zXa~!q=T30!LkCzMJ)Z`Sedq$qqh=4-TF&QtLLN0&fURY$7cY;Rr-Q9!Y$aaqKDe$k zzzuloInKD_jGYPg{){!B1y*y6wT!8GRz0_#C4KaCHrTTi{~mJ=SS`L+8Bdvl#Wx1FJtYlFNa@OGHrbzScTU~|;tZ-EQJ>T599 zyb7-tJzWI$T1L%_;pVI=_V5z0TEs5}yKmmV$20D7@C^KANzN5u_1O0-!H!dpUataM z+nDcqYh4YFcfEc2F1BZN^V8p<{0=z=^RqkVcgWrFdtmPSp7{88sB0MOJrn$fDt=?Z zUDI>qT?@BPJd;L&uR~M!?{%N)*MrsU&-(J%)3snVW6mXy*iB$HW6mp&-`s8n`@L8_ G*7`q|I}beo literal 0 HcmV?d00001 diff --git a/examples/Fluid2D/Shaders/SPIRV/clearGridIndices.comp.spv b/examples/Fluid2D/Shaders/SPIRV/clearGridIndices.comp.spv new file mode 100644 index 0000000000000000000000000000000000000000..d6051d5f640db176e44130ae1088c38a8718a9ca GIT binary patch literal 1096 zcmZ9LT}xX*5Qevj(WtG~T5G?wC#GMZNK3zw7KB0?@UGs6qetXGOp2!7_8;~4c~j!^ zob1xz!Z7>3^X|;->^b#XdoF}qVIefbt1wwBp#~GeVwkOb=iukT?!~CsegEN&iltDm zgld-2H1IRnSzZndY+`re%>rC8DhnKn zIdX?^dR`XYK~Wx`TpZ>{{S#w$V$9Vm@==age-D3C-+MdB{gl}M+|8=@(|)hhA0FlX z?%;BGoR7-kp!>bbH4`%Yulzhei+F1aceu?mcou77Kc@O=gIyi&Z;wt;y}g;km2ita zY#KKE;Ate;`7bubTHacVNKy~$}U@I2nL@}9qg&oF!N z{1RSH%?6dcqueHY?;>gdFJkW}w;l7YvUiE?^L~XU=hxV$UX%VI+;}|xv%@T{ z>AmGNYd*%C->z#ais;jAeLninZCx88fH&OYHACd%q3&f86jhmU@?W`L>@p qYmfd%-s3iYpB&EnLF9c0Yx*YZk>3Y7-zS*w=-U~y7dhV|)x8GL*+u;T literal 0 HcmV?d00001 diff --git a/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv b/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv new file mode 100644 index 0000000000000000000000000000000000000000..5d2a966a54d6e7115cd06e779efbfe4b75c77012 GIT binary patch literal 1920 zcmZvc+fP$L5XP6@te^-Y;0aBhi{)}E4_1??`59m$k%uxSZAlcZu3vVRFkixH1Y3z?9wPNJ2I^NqbBXIxk_ zY&W%ukycITYC%72H##jaIsjuGpZEK%Zg_=J_haia#Fx3TES6>&BbWepHzau6U=`!7kI>Waj zy)l`7B9TX6L2T6YfP2(JlLMMnVa={M|w^& zCRvpqdsf22CjVSy4fD-Qw=)_4zA*Bcy&#=@)L-;IdZst#KRjc~J@Y;=dc)=p$VKdP zuN!|`HhHNJADcS@vtAFqy{69L9ZA4Uw)8H@&)Z{hFTaKG|3gmhfp5oZVDoLnZN3fc zeev0MCG>HiY{X)lPk%oVhL1eh=G&JJhL4=syS~2z>AXGgKRV%##Civ3@piTqOOO1& zagH0>%gVt$@vVV#f7IkIdAs1oc_4k1!+QrW>vY`TBMEPYcY%-lB_DU1li;^}dFiA1 zmV|*@KE6lGhmX9L@3D0JKXeB2Kau1l;MCwvypZrtEROfF>hYS7Tk{zBPmjHnPCk6r zOF_Ez!fbfU+y^*(JJ&1eqvt9L12@k4e>ks&fm^)|=~fS4e3m8I_>H?6<5G+B(i=Ej Hdl!EJDh+Bm literal 0 HcmV?d00001 diff --git a/examples/Fluid2D/Shaders/SPIRV/render.frag.spv b/examples/Fluid2D/Shaders/SPIRV/render.frag.spv new file mode 100644 index 0000000000000000000000000000000000000000..275e7ec4d2af9bf99172dcc0553e7795a65b5227 GIT binary patch literal 2388 zcmZ9MYj4y=6vhX-U2c_&Tq=mL1r!uo=~4<Z**Hv7CN2oW5TQ-&V9;ZD`>lkpztwC;Yy2>+=p+vKM+ z|Kq=PXET2}XDt`8tzEd@+Ao1Ez#d&@wm#pkXZx7B1{~_;SAp?Z{~5N!{GT(M%ltEi zd++*!d+&1fUliLO7?=AFj7R=eZ1;<}?;_UUE3w#5?mJOGO>eye@86;CFFeD{=JG!M zW@4?}_dUoT6n+)G{@cLw%#+1l-T~&Y2iMp^t|s&iY|k6|yV&aOxq7)8_4dG@kLOqK zyIBG1eNXD;YC``tw*3hGH`wa!nR>Yz^*flq1ICScR(Zsl*!qllZuu4J@Z1l9|1kZ= z?eRnKFE?=+*lTra9DXyG(AiV}W%hkHvwipb@qZS5mLL5e9~ExLeM99ozjJ;akAg=f2_lr1-3(LDkm%DR8}d Date: Sat, 5 Aug 2023 19:45:50 +0900 Subject: [PATCH 03/13] fixup! add HLSL shaders --- .../Fluid2D/Shaders/HLSL_DX12/{buildVBID.comp => buildVBIB.comp} | 0 1 file changed, 0 insertions(+), 0 deletions(-) rename examples/Fluid2D/Shaders/HLSL_DX12/{buildVBID.comp => buildVBIB.comp} (100%) diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/buildVBID.comp b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp similarity index 100% rename from examples/Fluid2D/Shaders/HLSL_DX12/buildVBID.comp rename to examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp From 0b239c2a3349dee00c2e5aecda38d1c97e034b62 Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 5 Aug 2023 19:46:19 +0900 Subject: [PATCH 04/13] fixup! transpile shaders for other platforms --- .../GLSL_GL/{buildVBID.comp => buildVBIB.comp} | 0 .../GLSL_VULKAN/{buildVBID.comp => buildVBIB.comp} | 0 .../Metal/{buildVBID.comp => buildVBIB.comp} | 0 .../{buildVBID.comp.spv => buildVBIB.comp.spv} | Bin 4 files changed, 0 insertions(+), 0 deletions(-) rename examples/Fluid2D/Shaders/GLSL_GL/{buildVBID.comp => buildVBIB.comp} (100%) rename examples/Fluid2D/Shaders/GLSL_VULKAN/{buildVBID.comp => buildVBIB.comp} (100%) rename examples/Fluid2D/Shaders/Metal/{buildVBID.comp => buildVBIB.comp} (100%) rename examples/Fluid2D/Shaders/SPIRV/{buildVBID.comp.spv => buildVBIB.comp.spv} (100%) diff --git a/examples/Fluid2D/Shaders/GLSL_GL/buildVBID.comp b/examples/Fluid2D/Shaders/GLSL_GL/buildVBIB.comp similarity index 100% rename from examples/Fluid2D/Shaders/GLSL_GL/buildVBID.comp rename to examples/Fluid2D/Shaders/GLSL_GL/buildVBIB.comp diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBID.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp similarity index 100% rename from examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBID.comp rename to examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp diff --git a/examples/Fluid2D/Shaders/Metal/buildVBID.comp b/examples/Fluid2D/Shaders/Metal/buildVBIB.comp similarity index 100% rename from examples/Fluid2D/Shaders/Metal/buildVBID.comp rename to examples/Fluid2D/Shaders/Metal/buildVBIB.comp diff --git a/examples/Fluid2D/Shaders/SPIRV/buildVBID.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv similarity index 100% rename from examples/Fluid2D/Shaders/SPIRV/buildVBID.comp.spv rename to examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv From 04b8d540f629cfda53bbf7323c819bfc1c887db0 Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 5 Aug 2023 21:12:56 +0900 Subject: [PATCH 05/13] add render.vert --- .../Fluid2D/Shaders/HLSL_DX12/render.vert | 22 +++++++++++++++++++ 1 file changed, 22 insertions(+) create mode 100644 examples/Fluid2D/Shaders/HLSL_DX12/render.vert diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/render.vert b/examples/Fluid2D/Shaders/HLSL_DX12/render.vert new file mode 100644 index 0000000..4b01499 --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/render.vert @@ -0,0 +1,22 @@ +struct VS_INPUT{ + float3 Position : POSITION0; + float4 Color : COLOR0; + float2 UV1 : UV0; + float2 UV2 : UV1; +}; + +struct VS_OUTPUT{ + float4 Position : SV_POSITION; + float4 Color : COLOR0; + float2 UV1 : UV0; + float2 UV2 : UV1; +}; + +VS_OUTPUT main(VS_INPUT input){ + VS_OUTPUT output; + output.Position = float4(input.Position, 1.0f); + output.UV1 = input.UV1; + output.UV2 = input.UV2; + output.Color = input.Color; + return output; +} \ No newline at end of file From f867b5c3a7c213708a9b54e2bdc15f59d1be1caf Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 5 Aug 2023 21:13:15 +0900 Subject: [PATCH 06/13] transpile render.vert --- examples/Fluid2D/Shaders/GLSL_GL/render.vert | 56 +++++++++++++++ .../Fluid2D/Shaders/GLSL_VULKAN/render.vert | 53 ++++++++++++++ examples/Fluid2D/Shaders/Metal/render.vert | 67 ++++++++++++++++++ .../Fluid2D/Shaders/SPIRV/render.vert.spv | Bin 0 -> 2776 bytes 4 files changed, 176 insertions(+) create mode 100644 examples/Fluid2D/Shaders/GLSL_GL/render.vert create mode 100644 examples/Fluid2D/Shaders/GLSL_VULKAN/render.vert create mode 100644 examples/Fluid2D/Shaders/Metal/render.vert create mode 100644 examples/Fluid2D/Shaders/SPIRV/render.vert.spv diff --git a/examples/Fluid2D/Shaders/GLSL_GL/render.vert b/examples/Fluid2D/Shaders/GLSL_GL/render.vert new file mode 100644 index 0000000..831db67 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/render.vert @@ -0,0 +1,56 @@ +#version 430 + +out gl_PerVertex +{ + vec4 gl_Position; +}; + +struct VS_INPUT +{ + vec3 Position; + vec4 Color; + vec2 UV1; + vec2 UV2; +}; + +struct VS_OUTPUT +{ + vec4 Position; + vec4 Color; + vec2 UV1; + vec2 UV2; +}; + +layout(location = 0) in vec3 input_Position; +layout(location = 1) in vec4 input_Color; +layout(location = 2) in vec2 input_UV1; +layout(location = 3) in vec2 input_UV2; +layout(location = 0) out vec4 _entryPointOutput_Color; +layout(location = 1) out vec2 _entryPointOutput_UV1; +layout(location = 2) out vec2 _entryPointOutput_UV2; + +VS_OUTPUT _main(VS_INPUT _input) +{ + VS_OUTPUT _output; + _output.Position = vec4(_input.Position, 1.0); + _output.UV1 = _input.UV1; + _output.UV2 = _input.UV2; + _output.Color = _input.Color; + return _output; +} + +void main() +{ + VS_INPUT _input; + _input.Position = input_Position; + _input.Color = input_Color; + _input.UV1 = input_UV1; + _input.UV2 = input_UV2; + VS_INPUT param = _input; + VS_OUTPUT flattenTemp = _main(param); + gl_Position = flattenTemp.Position; + _entryPointOutput_Color = flattenTemp.Color; + _entryPointOutput_UV1 = flattenTemp.UV1; + _entryPointOutput_UV2 = flattenTemp.UV2; +} + diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/render.vert b/examples/Fluid2D/Shaders/GLSL_VULKAN/render.vert new file mode 100644 index 0000000..4bb6b7e --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/render.vert @@ -0,0 +1,53 @@ +#version 430 + +struct VS_INPUT +{ + vec3 Position; + vec4 Color; + vec2 UV1; + vec2 UV2; +}; + +struct VS_OUTPUT +{ + vec4 Position; + vec4 Color; + vec2 UV1; + vec2 UV2; +}; + +layout(location = 0) in vec3 input_Position; +layout(location = 1) in vec4 input_Color; +layout(location = 2) in vec2 input_UV1; +layout(location = 3) in vec2 input_UV2; +layout(location = 0) out vec4 _entryPointOutput_Color; +layout(location = 1) out vec2 _entryPointOutput_UV1; +layout(location = 2) out vec2 _entryPointOutput_UV2; + +VS_OUTPUT _main(VS_INPUT _input) +{ + VS_OUTPUT _output; + _output.Position = vec4(_input.Position, 1.0); + _output.UV1 = _input.UV1; + _output.UV2 = _input.UV2; + _output.Color = _input.Color; + return _output; +} + +void main() +{ + VS_INPUT _input; + _input.Position = input_Position; + _input.Color = input_Color; + _input.UV1 = input_UV1; + _input.UV2 = input_UV2; + VS_INPUT param = _input; + VS_OUTPUT flattenTemp = _main(param); + vec4 _position = flattenTemp.Position; + _position.y = -_position.y; + gl_Position = _position; + _entryPointOutput_Color = flattenTemp.Color; + _entryPointOutput_UV1 = flattenTemp.UV1; + _entryPointOutput_UV2 = flattenTemp.UV2; +} + diff --git a/examples/Fluid2D/Shaders/Metal/render.vert b/examples/Fluid2D/Shaders/Metal/render.vert new file mode 100644 index 0000000..a825159 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/render.vert @@ -0,0 +1,67 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct VS_INPUT +{ + float3 Position; + float4 Color; + float2 UV1; + float2 UV2; +}; + +struct VS_OUTPUT +{ + float4 Position; + float4 Color; + float2 UV1; + float2 UV2; +}; + +struct main0_out +{ + float4 _entryPointOutput_Color [[user(locn0)]]; + float2 _entryPointOutput_UV1 [[user(locn1)]]; + float2 _entryPointOutput_UV2 [[user(locn2)]]; + float4 gl_Position [[position]]; +}; + +struct main0_in +{ + float3 input_Position [[attribute(0)]]; + float4 input_Color [[attribute(1)]]; + float2 input_UV1 [[attribute(2)]]; + float2 input_UV2 [[attribute(3)]]; +}; + +static inline __attribute__((always_inline)) +VS_OUTPUT _main(thread const VS_INPUT& _input) +{ + VS_OUTPUT _output; + _output.Position = float4(_input.Position, 1.0); + _output.UV1 = _input.UV1; + _output.UV2 = _input.UV2; + _output.Color = _input.Color; + return _output; +} + +vertex main0_out main0(main0_in in [[stage_in]]) +{ + main0_out out = {}; + VS_INPUT _input; + _input.Position = in.input_Position; + _input.Color = in.input_Color; + _input.UV1 = in.input_UV1; + _input.UV2 = in.input_UV2; + VS_INPUT param = _input; + VS_OUTPUT flattenTemp = _main(param); + out.gl_Position = flattenTemp.Position; + out._entryPointOutput_Color = flattenTemp.Color; + out._entryPointOutput_UV1 = flattenTemp.UV1; + out._entryPointOutput_UV2 = flattenTemp.UV2; + return out; +} + diff --git a/examples/Fluid2D/Shaders/SPIRV/render.vert.spv b/examples/Fluid2D/Shaders/SPIRV/render.vert.spv new file mode 100644 index 0000000000000000000000000000000000000000..fe8e5d88b165633c37ff3492c7dcb374dde1cc4e GIT binary patch literal 2776 zcmZvdSyNL%5QPUw0t)U6s0hIo7gTl-6#*4gl&~bO7*iOmR6&vo7fXH7U*zxeMa%Qe zorZ{BtC>Eh&-Cr-=^L9G&o?{Q;P$%~_uj?N5!WcjxmK6*>+JmU{Ee;X!;PWgTPO~? zCND&D$Q|&0LVq_4r81yjHY^*L-H|QHauO)vr*RP#1rs8aFKs%DIO}_FkWLM%B zkROg`Z+`f_`QhsGgLm_{sVC0mX-4l>RNG#Q`s^@$pVkLuL$U#^e?rwH)KBS4JXzxh`@<@@X9`Q2(c7d?;pxEpZXLCnz&@+bYB7&qX#gqze_Ez}B|K{KUw zeWMUX#qvsVv#MPQH@!#mP&i-Ry^y5(m=X7JBcCbOa>ZIy{46Lo=JAQyj+ATmdiv}6 zQaM^KeJ!H17=2c3W}{SnP}+(L<+Y+}Y4S6LE~&X?w(`YtRQr;tD8+*BT;6)n;^*$) zxb>sO&*cB&^k!D)rf*#`9x#0dQ#Y6%gQ*qF%!8SUIaNqccv|CShhsk+tLexA$KF}} zWv}Jsob`3^J=NF2w`Fx4(=PV)CI~89DKANsDXVgNR7-v}2%K4Zh-U(g` z9PyRl+Zn*j8GDGo3cj5I%-j(tA2|L^cWy*A&}-(02hC40Nv&wPcYe#PpJ0A>3^NZ? zb(r{c9p--d%|U-x`B?czxsxfqh4~ni@s8uNd#V7820dpSk#R@p$u}xP-z!58MuVO+ z-&w44ycBgR^rz@_yV0@^SWA89aLW2u9CbJ(a;ROX!)Qc^Uk-fS(1<_B&$V zvG*;9IC|zPD`WQTZop`)<`>?N@4|At6bDa!<`RrPt@Vxf9d{;hjaoCx#C>+k*pHrghs=7}C7*h%*I@M4iyryZ)-AJM zoRbf4n)C9>frkEo(OCSVd~z_`mgACqc+*^#PY#=ZFdE`~e^+GW;QJ#7=ejC`H_bKq WzjYt@+}$7Sui1eB literal 0 HcmV?d00001 From 637e13fe7b2acead5b52dc4b2efa537427e80afb Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 5 Aug 2023 21:13:59 +0900 Subject: [PATCH 07/13] implementing Fluid2D --- examples/CMakeLists.txt | 1 + examples/Fluid2D/CMakeLists.txt | 61 ++++ examples/Fluid2D/Fluid2D.cpp | 601 ++++++++++++++++++++++++++++++++ examples/Fluid2D/Fluid2D.h | 145 ++++++++ examples/Fluid2D/Main.cpp | 59 ++++ 5 files changed, 867 insertions(+) create mode 100644 examples/Fluid2D/CMakeLists.txt create mode 100644 examples/Fluid2D/Fluid2D.cpp create mode 100644 examples/Fluid2D/Fluid2D.h create mode 100644 examples/Fluid2D/Main.cpp diff --git a/examples/CMakeLists.txt b/examples/CMakeLists.txt index 2516de1..360c050 100644 --- a/examples/CMakeLists.txt +++ b/examples/CMakeLists.txt @@ -4,3 +4,4 @@ add_subdirectory("ImGuiPlatform") add_subdirectory("glfw") add_subdirectory("imgui") add_subdirectory("GPUParticle") +add_subdirectory("Fluid2D") diff --git a/examples/Fluid2D/CMakeLists.txt b/examples/Fluid2D/CMakeLists.txt new file mode 100644 index 0000000..c736926 --- /dev/null +++ b/examples/Fluid2D/CMakeLists.txt @@ -0,0 +1,61 @@ +add_executable( + example_Fluid2D + Fluid2D.h + Fluid2D.cpp + Main.cpp +) + +target_include_directories( + example_Fluid2D + PRIVATE + ../../src/ +) + +target_link_libraries( + example_Fluid2D + PRIVATE + LLGI + glfw +) + +target_link_directories(example_Fluid2D PRIVATE ${THIRDPARTY_LIBRARY_DIRECTORIES}) +target_link_libraries(example_Fluid2D PRIVATE ${THIRDPARTY_LIBRARIES}) +target_compile_definitions(example_Fluid2D PRIVATE EXAMPLE_ASSET_DIR="${CMAKE_CURRENT_LIST_DIR}") + +if(WIN32) + # None +elseif(APPLE) + + find_library(COCOA_LIBRARY Cocoa) + find_library(METAL_LIBRARY Metal) + find_library(APPKIT_LIBRARY AppKit) + find_library(METALKIT_LIBRARY MetalKit) + find_library(QUARTZ_CORE_LIBRARY QuartzCore) + + set(EXTRA_LIBS ${COCOA_LIBRARY} ${APPKIT_LIBRARY} ${METAL_LIBRARY} ${METALKIT_LIBRARY} ${QUARTZ_CORE_LIBRARY}) + target_link_libraries(example_Fluid2D PRIVATE ${EXTRA_LIBS}) + +else() + + find_package(Threads REQUIRED) + target_link_libraries( + example_Fluid2D + PRIVATE + ${CMAKE_THREAD_LIBS_INIT} + pthread + X11 + X11-xcb + ) + +endif() + + +if(BUILD_VULKAN) + find_package(Vulkan REQUIRED) + target_include_directories(example_Fluid2D PRIVATE ${Vulkan_INCLUDE_DIRS}) + target_link_libraries(example_Fluid2D PRIVATE ${Vulkan_LIBRARIES}) + + target_include_directories(example_Fluid2D PRIVATE ${LLGI_THIRDPARTY_INCLUDES}) + target_link_libraries(example_Fluid2D PRIVATE ${LLGI_THIRDPARTY_LIBRARIES}) + target_link_directories(example_Fluid2D PRIVATE ${LLGI_THIRDPARTY_LIBRARY_DIRECTORIES}) +endif() diff --git a/examples/Fluid2D/Fluid2D.cpp b/examples/Fluid2D/Fluid2D.cpp new file mode 100644 index 0000000..de0ed4b --- /dev/null +++ b/examples/Fluid2D/Fluid2D.cpp @@ -0,0 +1,601 @@ +#include + +#include +#include +#include +#include +#include +#include +#include + +#include "Fluid2D.h" + +std::vector LoadData(const char* path) +{ + std::vector ret; + +#ifdef _WIN32 + FILE* fp = nullptr; + fopen_s(&fp, path, "rb"); + +#else + FILE* fp = fopen(path, "rb"); +#endif + + if (fp == nullptr) + { + std::cout << "Not found : " << path << std::endl; + return ret; + } + + fseek(fp, 0, SEEK_END); + auto size = ftell(fp); + fseek(fp, 0, SEEK_SET); + + ret.resize(size); + fread(ret.data(), 1, size, fp); + fclose(fp); + + return ret; +} + +std::shared_ptr CreateShader( + LLGI::Graphics* graphics, + LLGI::DeviceType deviceType, + const char* csFilename, + LLGI::ShaderStageType stageType) +{ + const auto compiler = LLGI::CreateSharedPtr(LLGI::CreateCompiler(deviceType)); + + std::vector data; + + auto csBinaryPath = std::string(EXAMPLE_ASSET_DIR); + + if (compiler == nullptr) + { + csBinaryPath = csBinaryPath + "/Shaders/SPIRV/" + csFilename + ".spv"; + const auto binary_cs = LoadData(csBinaryPath.c_str()); + + LLGI::DataStructure d; + + d.Data = binary_cs.data(); + d.Size = static_cast(binary_cs.size()); + + data.push_back(d); + + return LLGI::CreateSharedPtr(graphics->CreateShader(data.data(), static_cast(data.size()))); + } + else + { + csBinaryPath = csBinaryPath + "/Shaders/HLSL_DX12/" + csFilename; + LLGI::CompilerResult result; + + auto code = LoadData(csBinaryPath.c_str()); + code.push_back(0); + + compiler->Compile(result, (const char*)code.data(), stageType); + + std::cout << result.Message.c_str() << std::endl; + + for (auto& b : result.Binary) + { + LLGI::DataStructure d; + d.Data = b.data(); + d.Size = static_cast(b.size()); + data.push_back(d); + } + + return LLGI::CreateSharedPtr(graphics->CreateShader(data.data(), static_cast(data.size()))); + } +} + +float Fluid2D::CalcRestDensity(const float h) +{ + const auto a = 4.0f / (Pi * std::powf(h, 8.0f)); + const auto l = 2 * ParticleRadius; + const auto n = (int)std::ceilf(h / l) + 1; + + auto r0 = 0.0f; + + for (int x = -n; x <= n; ++x) + { + for (int y = -n; y <= n; ++y) + { + const auto rij = LLGI::Vec2F(x * l, y * l); + const auto r = std::sqrtf(rij.X * rij.X + rij.Y * rij.Y); + if (r >= 0.0 && r <= h) + { + const auto q = h * h - r * r; + r0 += a * q * q * q; + } + } + } + return r0; +} + +void copyVec2F(const LLGI::Vec2F& src, float* dst) +{ + dst[0] = src.X; + dst[1] = src.Y; +} + +void copyVec2I(const LLGI::Vec2I& src, int* dst) +{ + dst[0] = src.X; + dst[1] = src.Y; +} + +Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) + : gridNum_(LLGI::Vec2I(60, 60)) + , gridSize_(LLGI::Vec2I(5, 5)) + , density_(CalcRestDensity(EffectiveRadius)) + , wpoly6_(4.0f / (Pi * std::powf(EffectiveRadius, 8.0f))) + , gwspiky_(-30.0f / (Pi * powf(EffectiveRadius, 5.0f))) +{ + std::cout << "Creating compute buffers." << std::endl; + + particleComputeIndex_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Compute | LLGI::BufferUsageType::CopySrc, + sizeof(int) * ParticlesCount * 6)); + + particleComputeVertex_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Compute | LLGI::BufferUsageType::CopySrc, + sizeof(Vertex) * ParticlesCount * 4)); + + particleIndex_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Index, + sizeof(int) * ParticlesCount * 6)); + + particleVertex_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Vertex, + sizeof(Vertex) * ParticlesCount * 4)); + + particles_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Compute | LLGI::BufferUsageType::CopyDst, + sizeof(Particle) * ParticlesCount)); + + gridTable_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Compute, + sizeof(int[2]) * ParticlesCount)); + + gridIndicesTable_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer( + LLGI::BufferUsageType::Compute, + sizeof(int[2]) * gridNum_.X * gridNum_.Y)); + + // Create Shader + std::cout << "Creating shaders." << std::endl; + + vs_ = CreateShader(graphics, deviceType, "render.vert", LLGI::ShaderStageType::Vertex); + ps_ = CreateShader(graphics, deviceType, "render.frag", LLGI::ShaderStageType::Pixel); + calcExternalShader_ = CreateShader(graphics, deviceType, "calcExternalForce.comp", LLGI::ShaderStageType::Compute); + buildGridShader_ = CreateShader(graphics, deviceType, "buildGrid.comp", LLGI::ShaderStageType::Compute); + bitonicSortShader_ = CreateShader(graphics, deviceType, "bitonicSort.comp", LLGI::ShaderStageType::Compute); + clearGridIndicesShader_ = CreateShader(graphics, deviceType, "clearGridIndices.comp", LLGI::ShaderStageType::Compute); + buildGridIndicesShader_ = CreateShader(graphics, deviceType, "buildGridIndices.comp", LLGI::ShaderStageType::Compute); + calcScalingFactorShader_ = CreateShader(graphics, deviceType, "calcScalingFactor.comp", LLGI::ShaderStageType::Compute); + calcCorrectPositionShader_ = CreateShader(graphics, deviceType, "calcCorrectPosition.comp", LLGI::ShaderStageType::Compute); + integrateShader_ = CreateShader(graphics, deviceType, "integrate.comp", LLGI::ShaderStageType::Compute); + buildVBIBShader_ = CreateShader(graphics, deviceType, "buildVBIB.comp", LLGI::ShaderStageType::Compute); + + // calcExternal + std::cout << "Creating calcExternalPipeline_" << std::endl; + calcExternalPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + calcExternalConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(CalcExternalForceCB)) + ); + calcExternalPipeline_->SetShader(LLGI::ShaderStageType::Compute, calcExternalShader_.get()); + if (!calcExternalPipeline_->Compile()) + { + std::cerr << "failed to compile calcExternalPipline" << std::endl; + abort(); + } + + { + const auto cb = (CalcExternalForceCB*)calcExternalConstant_->Lock(); + + cb->Force[0] = 0; + cb->Force[1] = 0; + cb->Gravity[0] = 0; + cb->Gravity[1] = 100; + cb->Dt[0] = Dt; + cb->Dt[1] = 0; + copyVec2I(gridNum_, cb->GridNum); + copyVec2I(gridSize_, cb->GridSize); + + calcExternalConstant_->Unlock(); + } + + // buildGrid + std::cout << "Creating buildGridPipeline_" << std::endl; + buildGridPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + buildGridConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(BuildGridCB)) + ); + + buildGridPipeline_->SetShader(LLGI::ShaderStageType::Compute, buildGridShader_.get()); + if (!buildGridPipeline_->Compile()) + { + std::cerr << "failed to compile buildGridPipeline" << std::endl; + abort(); + } + + { + const auto cb = (BuildGridCB*)buildGridConstant_->Lock(); + + copyVec2I(gridNum_, cb->GridNum); + copyVec2I(gridSize_, cb->GridSize); + cb->ParticlesCount = ParticlesCount; + + buildGridConstant_->Unlock(); + } + + // bitonicSort + std::cout << "Creating bitonicSortPipeline_" << std::endl; + bitonicSortPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + bitonicSortConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(BuildGridIndicesCB)) + ); + + bitonicSortPipeline_->SetShader(LLGI::ShaderStageType::Compute, bitonicSortShader_.get()); + if (!bitonicSortPipeline_->Compile()) + { + std::cerr << "failed to compile bitonicSortPipeline" << std::endl; + abort(); + } + + // clearGridIndices + std::cout << "Creating clearGridIndicesPipeline_" << std::endl; + clearGridIndicesPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + clearGridIndicesPipeline_->SetShader(LLGI::ShaderStageType::Compute, clearGridIndicesShader_.get()); + if (!clearGridIndicesPipeline_->Compile()) + { + std::cerr << "failed to compile clearGridIndicesPipeline" << std::endl; + abort(); + } + + // buildGridIndices + std::cout << "Creating buildGridIndicesPipeline_" << std::endl; + buildGridIndicesPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + buildGridIndicesConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(BuildGridIndicesCB)) + ); + buildGridIndicesPipeline_->SetShader(LLGI::ShaderStageType::Compute, buildGridIndicesShader_.get()); + if (!buildGridIndicesPipeline_->Compile()) + { + std::cerr << "failed to compile buildGridIndicesPipeline" << std::endl; + abort(); + } + + { + const auto cb = (BuildGridIndicesCB*)buildGridIndicesConstant_->Lock(); + + cb->ParticlesCount = ParticlesCount; + + buildGridIndicesConstant_->Unlock(); + } + + // calcScalingfactor + std::cout << "Creating calcScalingFactorPipeline_" << std::endl; + calcScalingFactorPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + calcScalingFactorConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(CalcScalingFactorCB)) + ); + calcScalingFactorPipeline_->SetShader(LLGI::ShaderStageType::Compute, calcScalingFactorShader_.get()); + if (!calcScalingFactorPipeline_->Compile()) + { + std::cerr << "failed to compile calcScalingFactorPipeline" << std::endl; + abort(); + } + + { + const auto cb = (CalcScalingFactorCB*)calcScalingFactorConstant_->Lock(); + + copyVec2I(gridNum_, cb->GridNum); + copyVec2I(gridSize_, cb->GridSize); + cb->EffectiveRadius = EffectiveRadius; + cb->Density = density_; + cb->Eps = Eps; + cb->Dt = Dt; + cb->Wpoly6 = wpoly6_; + cb->GWspiky = gwspiky_; + + calcScalingFactorConstant_->Unlock(); + } + + // calcCorrectPosition + std::cout << "Creating calcCorrectPositionPipeline_" << std::endl; + calcCorrectPositionPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + calcCorrectPositionConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(CalcCorrectPositionCB)) + ); + + calcCorrectPositionPipeline_->SetShader(LLGI::ShaderStageType::Compute, calcCorrectPositionShader_.get()); + if (!calcCorrectPositionPipeline_->Compile()) + { + std::cerr << "failed to compile calcCorrectPositionPipeline" << std::endl; + abort(); + } + + { + const auto cb = (CalcCorrectPositionCB*)calcCorrectPositionConstant_->Lock(); + + copyVec2I(gridNum_, cb->GridNum); + copyVec2I(gridSize_, cb->GridSize); + cb->EffectiveRadius = EffectiveRadius; + cb->Density = density_; + cb->Eps = Eps; + cb->Dt = Dt; + cb->Wpoly6 = wpoly6_; + cb->GWspiky = gwspiky_; + + calcCorrectPositionConstant_->Unlock(); + } + + // integrate + std::cout << "Creating integratePipeline_" << std::endl; + integratePipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + integrateConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(IntegrateCB)) + ); + + integratePipeline_->SetShader(LLGI::ShaderStageType::Compute, integrateShader_.get()); + if (!integratePipeline_->Compile()) + { + std::cerr << "failed to compile integratePipeline" << std::endl; + abort(); + } + + { + const auto cb = (IntegrateCB*)integrateConstant_->Lock(); + + cb->Dt = Dt; + + integrateConstant_->Unlock(); + } + + // buildVBIB + std::cout << "Creating buildVBIBPipeline_" << std::endl; + buildVBIBPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + buildVBIBConstant_ = LLGI::CreateSharedPtr( + graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(BuildVBIBCB)) + ); + + buildVBIBPipeline_->SetShader(LLGI::ShaderStageType::Compute, buildVBIBShader_.get()); + if (!buildVBIBPipeline_->Compile()) + { + std::cerr << "failed to compile buildVBIBPipeline" << std::endl; + abort(); + } + + { + const auto cb = (BuildVBIBCB*)buildVBIBConstant_->Lock(); + + cb->ParticleRadius = ParticleRadius; + cb->Color[0] = 0.2f; + cb->Color[1] = 0.2f; + cb->Color[2] = 1.0f; + cb->Color[3] = 1.0f; + + cb->FixedColor[0] = 0.0f; + cb->FixedColor[1] = 0.0f; + cb->FixedColor[2] = 0.0f; + cb->FixedColor[3] = 1.0f; + + buildVBIBConstant_->Unlock(); + } +} + +void Fluid2D::Initialize(LLGI::Graphics* graphics, LLGI::CommandList* commandList) +{ + std::cout << "Initializing" << std::endl; + + const auto particlesInput = LLGI::CreateUniqueReference( + graphics->CreateBuffer( + LLGI::BufferUsageType::MapWrite | LLGI::BufferUsageType::CopySrc, + sizeof(Particle) * ParticlesCount + ) + ); + + { + const auto data = (Particle*)particlesInput->Lock(); + + int ind = 0; + for (int y = 0; y < 300 / (ParticleRadius * 2); y++) + { + for (int x = 0; x < 300 / (ParticleRadius * 2); x++) + { + if (x * (ParticleRadius * 2) + ParticleRadius < 8 || + x * (ParticleRadius * 2) + ParticleRadius > 300 - 8 || + y * (ParticleRadius * 2) + ParticleRadius < 8 || + y * (ParticleRadius * 2) + ParticleRadius > 300 - 8) + { + const auto posX = x * (ParticleRadius * 2) + ParticleRadius; + const auto posY = y * (ParticleRadius * 2) + ParticleRadius; + data[ind].Current[0] = posX; + data[ind].Current[1] = posY; + data[ind].Next[0] = posX; + data[ind].Next[1] = posY; + data[ind].Velocity[0] = 0; + data[ind].Velocity[1] = 0; + data[ind].Pscl = 0; + data[ind].IsFix = true; + ind++; + } + } + } + + for (int i = ind; i < ParticlesCount; i++) + { + data[i].Current[0] = static_cast(i % 50 * 4 + 10); + data[i].Current[1] = static_cast(i / 50 * 4 + 20); + data[i].Velocity[0] = 0; + data[i].Velocity[1] = 0; + data[i].Pscl = 0; + } + + particlesInput->Unlock(); + } + + commandList->Begin(); + commandList->CopyBuffer(particlesInput.get(), particles_.get()); + commandList->End(); + graphics->Execute(commandList); + graphics->WaitFinish(); +} + +void Fluid2D::Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList) +{ + commandList->Begin(); + + commandList->BeginComputePass(); + + commandList->SetPipelineState(calcExternalPipeline_.get()); + commandList->SetConstantBuffer(calcExternalConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + commandList->ResetComputeBuffer(); + + for (int l = 0; l < IterationCount; l++) + { + commandList->SetPipelineState(buildGridPipeline_.get()); + commandList->SetConstantBuffer(buildGridConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1, LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + commandList->ResetComputeBuffer(); + + const int nlog = static_cast(std::log2f(ParticlesCount)); + int inc; + + commandList->SetPipelineState(bitonicSortPipeline_.get()); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0, LLGI::ShaderStageType::Compute); + for (int i = 0; i < nlog; i++) + { + inc = 1 << i; + + const auto dir = 2 << i; + + for (int j = 0; j < i + 1; j++) + { + { + const auto cb = (BitonicSortCB*)bitonicSortConstant_->Lock(); + cb->Inc = inc; + cb->Dir = dir; + bitonicSortConstant_->Unlock(); + } + + commandList->SetConstantBuffer(bitonicSortConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount / 2, 1, 1, 1, 1, 1); + + inc /= 2; + } + } + commandList->ResetComputeBuffer(); + + commandList->SetPipelineState(clearGridIndicesPipeline_.get()); + commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 0, LLGI::ShaderStageType::Compute); + commandList->Dispatch(gridNum_.X * gridNum_.Y, 1, 1, 1, 1, 1); + commandList->ResetComputeBuffer(); + + commandList->SetPipelineState(buildGridIndicesPipeline_.get()); + commandList->SetConstantBuffer(buildGridIndicesConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 1, LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + commandList->ResetComputeBuffer(); + + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 2, LLGI::ShaderStageType::Compute); + { + commandList->SetPipelineState(calcScalingFactorPipeline_.get()); + commandList->SetConstantBuffer(calcScalingFactorConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + + commandList->SetPipelineState(calcCorrectPositionPipeline_.get()); + commandList->SetConstantBuffer(calcCorrectPositionConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + } + commandList->ResetComputeBuffer(); + } + + commandList->SetPipelineState(integratePipeline_.get()); + commandList->SetConstantBuffer(integrateConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + commandList->ResetComputeBuffer(); + + commandList->SetPipelineState(buildVBIBPipeline_.get()); + commandList->SetConstantBuffer(buildVBIBConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particleComputeVertex_.get(), sizeof(Vertex), 1, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particleComputeIndex_.get(), sizeof(int), 2, LLGI::ShaderStageType::Compute); + commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + commandList->ResetComputeBuffer(); + + commandList->EndComputePass(); + + commandList->CopyBuffer(particleComputeVertex_.get(), particleVertex_.get()); + commandList->CopyBuffer(particleComputeIndex_.get(), particleIndex_.get()); + + commandList->End(); + graphics->Execute(commandList); + graphics->WaitFinish(); +} + +void Fluid2D::Render(LLGI::Graphics* graphics, LLGI::CommandList* commandList, LLGI::RenderPass* renderPass) +{ + const auto renderPassState = LLGI::CreateSharedPtr(graphics->CreateRenderPassPipelineState(renderPass)); + + std::shared_ptr pipeline; + auto itr = pipelineCache_.find(renderPassState); + if (itr == pipelineCache_.end()) + { + const auto pip = graphics->CreatePiplineState(); + pip->VertexLayouts[0] = LLGI::VertexLayoutFormat::R32G32B32_FLOAT; + pip->VertexLayouts[1] = LLGI::VertexLayoutFormat::R8G8B8A8_UNORM; + pip->VertexLayouts[3] = LLGI::VertexLayoutFormat::R32G32_FLOAT; + pip->VertexLayouts[4] = LLGI::VertexLayoutFormat::R32G32_FLOAT; + pip->VertexLayoutNames[0] = "POSITION"; + pip->VertexLayoutNames[1] = "COLOR"; + pip->VertexLayoutNames[2] = "UV1"; + pip->VertexLayoutNames[3] = "UV2"; + pip->VertexLayoutCount = 4; + pip->Topology = LLGI::TopologyType::Triangle; + + pip->IsBlendEnabled = true; + pip->BlendSrcFunc = LLGI::BlendFuncType::SrcAlpha; + pip->BlendDstFunc = LLGI::BlendFuncType::One; + pip->BlendSrcFuncAlpha = LLGI::BlendFuncType::One; + pip->BlendDstFuncAlpha = LLGI::BlendFuncType::One; + pip->BlendEquationRGB = LLGI::BlendEquationType::Max; + pip->BlendEquationAlpha = LLGI::BlendEquationType::Max; + + pip->SetShader(LLGI::ShaderStageType::Vertex, vs_.get()); + pip->SetShader(LLGI::ShaderStageType::Pixel, ps_.get()); + pip->SetRenderPassPipelineState(renderPassState.get()); + pip->Compile(); + pipeline = LLGI::CreateSharedPtr(pip); + pipelineCache_[renderPassState] = pipeline; + } + else + { + pipeline = itr->second; + } + + commandList->BeginRenderPass(renderPass); + + commandList->SetVertexBuffer(particleVertex_.get(), sizeof(Vertex), 0); + commandList->SetIndexBuffer(particleIndex_.get(), sizeof(int)); + commandList->SetPipelineState(pipeline.get()); + commandList->Draw(ParticlesCount * 2); + + commandList->EndRenderPass(); +} \ No newline at end of file diff --git a/examples/Fluid2D/Fluid2D.h b/examples/Fluid2D/Fluid2D.h new file mode 100644 index 0000000..9487d65 --- /dev/null +++ b/examples/Fluid2D/Fluid2D.h @@ -0,0 +1,145 @@ +#pragma once +#include +#include +#include +#include + +class Fluid2D; + +struct alignas(16) CalcExternalForceCB +{ + float Force[2]; + float Gravity[2]; + float Dt[2]; + int GridNum[2]; + int GridSize[2]; +}; + +struct alignas(16) BuildGridCB +{ + int GridNum[2]; + int GridSize[2]; + int ParticlesCount; +}; + +struct alignas(16) BitonicSortCB +{ + int Inc; + int Dir; +}; + +struct alignas(16) BuildGridIndicesCB +{ + int ParticlesCount; +}; + +struct alignas(16) CalcScalingFactorCB +{ + int GridNum[2]; + int GridSize[2]; + float EffectiveRadius; + float Density; + float Eps; + float Dt; + float Wpoly6; + float GWspiky; +}; + +using CalcCorrectPositionCB = CalcScalingFactorCB; + +struct alignas(16) IntegrateCB +{ + float Dt; +}; + +struct alignas(16) BuildVBIBCB +{ + float ParticleRadius; + float Color[4]; + float FixedColor[4]; +}; + +struct alignas(16) Vertex +{ + float Position[3]; + float Color[3]; + float UV1[2]; + float UV2[2]; +}; + +struct alignas(16) Particle +{ + float Current[2]; + float Next[2]; + float Velocity[2]; + float Pscl; + bool IsFix; +}; + +class Fluid2D +{ +public: + Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType); + void Initialize(LLGI::Graphics* graphics, LLGI::CommandList* commandList); + void Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList); + void Render(LLGI::Graphics* graphics, LLGI::CommandList* commandList, LLGI::RenderPass* renderPass); + +private: + std::shared_ptr vs_; + std::shared_ptr ps_; + + std::shared_ptr particleComputeIndex_; + std::shared_ptr particleComputeVertex_; + std::shared_ptr particleIndex_; + std::shared_ptr particleVertex_; + std::shared_ptr particles_; + std::shared_ptr gridTable_; + std::shared_ptr gridIndicesTable_; + + std::shared_ptr calcExternalShader_; + std::shared_ptr buildGridShader_; + std::shared_ptr bitonicSortShader_; + std::shared_ptr clearGridIndicesShader_; + std::shared_ptr buildGridIndicesShader_; + std::shared_ptr calcScalingFactorShader_; + std::shared_ptr calcCorrectPositionShader_; + std::shared_ptr integrateShader_; + std::shared_ptr buildVBIBShader_; + + std::shared_ptr calcExternalConstant_; + std::shared_ptr buildGridConstant_; + std::shared_ptr bitonicSortConstant_; + std::shared_ptr buildGridIndicesConstant_; + std::shared_ptr calcScalingFactorConstant_; + std::shared_ptr calcCorrectPositionConstant_; + std::shared_ptr integrateConstant_; + std::shared_ptr buildVBIBConstant_; + + std::shared_ptr calcExternalPipeline_; + std::shared_ptr buildGridPipeline_; + std::shared_ptr bitonicSortPipeline_; + std::shared_ptr clearGridIndicesPipeline_; + std::shared_ptr buildGridIndicesPipeline_; + std::shared_ptr calcScalingFactorPipeline_; + std::shared_ptr calcCorrectPositionPipeline_; + std::shared_ptr integratePipeline_; + std::shared_ptr buildVBIBPipeline_; + + std::unordered_map, std::shared_ptr> pipelineCache_; + + LLGI::Vec2I gridNum_; + LLGI::Vec2I gridSize_; + float density_; + float wpoly6_; + float gwspiky_; + + static float CalcRestDensity(float h); + + static constexpr float Pi = 3.14159265f; + static constexpr int IterationCount = 4; + static constexpr float EffectiveRadius = 6.0f; + static constexpr float ParticleRadius = 2.0f; + static constexpr int ParticlesCount = 2048; + static constexpr float Eps = 0.1f; + static constexpr float Dt = 1.0f / 60.0f; +}; \ No newline at end of file diff --git a/examples/Fluid2D/Main.cpp b/examples/Fluid2D/Main.cpp new file mode 100644 index 0000000..040b9f3 --- /dev/null +++ b/examples/Fluid2D/Main.cpp @@ -0,0 +1,59 @@ +#include +#include +#include +#include +#include +#include "Fluid2D.h" + +#ifdef _WIN32 +#pragma comment(lib, "d3dcompiler.lib") +#endif + +int main() +{ + LLGI::PlatformParameter pp; + pp.Device = LLGI::DeviceType::Default; + pp.WaitVSync = true; + auto window = std::unique_ptr(LLGI::CreateWindow("Fluid2D", LLGI::Vec2I(1280, 720))); + auto platform = LLGI::CreatePlatform(pp, window.get()); + + auto graphics = platform->CreateGraphics(); + auto sfMemoryPool = graphics->CreateSingleFrameMemoryPool(1024 * 1024, 128); + auto commandList = graphics->CreateCommandList(sfMemoryPool); + + const auto fluid2d = std::make_unique(graphics, pp.Device); + fluid2d->Initialize(graphics, commandList); + + bool initialized = false; + + while (true) + { + if (!platform->NewFrame()) break; + + sfMemoryPool->NewFrame(); + + fluid2d->Update(graphics, commandList); + + const LLGI::Color8 color = { 50, 50, 50, 255 }; + auto renderPass = platform->GetCurrentScreen(color, true, false); + + commandList->WaitUntilCompleted(); + + commandList->Begin(); + fluid2d->Render(graphics, commandList, renderPass); + commandList->End(); + graphics->Execute(commandList); + graphics->WaitFinish(); + + platform->Present(); + } + + graphics->WaitFinish(); + + LLGI::SafeRelease(sfMemoryPool); + LLGI::SafeRelease(commandList); + LLGI::SafeRelease(graphics); + LLGI::SafeRelease(platform); + + return 0; +} \ No newline at end of file From 49db265f210d7e1aee6e2725fd7a3f511d7a8a36 Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 19 Aug 2023 09:01:06 +0900 Subject: [PATCH 08/13] fix --- examples/Fluid2D/Fluid2D.cpp | 44 ++++++++++++++++++------------------ 1 file changed, 22 insertions(+), 22 deletions(-) diff --git a/examples/Fluid2D/Fluid2D.cpp b/examples/Fluid2D/Fluid2D.cpp index de0ed4b..08971f2 100644 --- a/examples/Fluid2D/Fluid2D.cpp +++ b/examples/Fluid2D/Fluid2D.cpp @@ -457,17 +457,17 @@ void Fluid2D::Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList) commandList->BeginComputePass(); commandList->SetPipelineState(calcExternalPipeline_.get()); - commandList->SetConstantBuffer(calcExternalConstant_.get(), LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(calcExternalConstant_.get(), 0); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); commandList->ResetComputeBuffer(); for (int l = 0; l < IterationCount; l++) { commandList->SetPipelineState(buildGridPipeline_.get()); - commandList->SetConstantBuffer(buildGridConstant_.get(), LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1, LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(buildGridConstant_.get(), 0); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); commandList->ResetComputeBuffer(); @@ -475,7 +475,7 @@ void Fluid2D::Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList) int inc; commandList->SetPipelineState(bitonicSortPipeline_.get()); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0); for (int i = 0; i < nlog; i++) { inc = 1 << i; @@ -491,7 +491,7 @@ void Fluid2D::Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList) bitonicSortConstant_->Unlock(); } - commandList->SetConstantBuffer(bitonicSortConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(bitonicSortConstant_.get(), 0); commandList->Dispatch(ParticlesCount / 2, 1, 1, 1, 1, 1); inc /= 2; @@ -500,43 +500,43 @@ void Fluid2D::Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList) commandList->ResetComputeBuffer(); commandList->SetPipelineState(clearGridIndicesPipeline_.get()); - commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 0, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 0); commandList->Dispatch(gridNum_.X * gridNum_.Y, 1, 1, 1, 1, 1); commandList->ResetComputeBuffer(); commandList->SetPipelineState(buildGridIndicesPipeline_.get()); - commandList->SetConstantBuffer(buildGridIndicesConstant_.get(), LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0, LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 1, LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(buildGridIndicesConstant_.get(), 0); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0); + commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 1); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); commandList->ResetComputeBuffer(); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1, LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 2, LLGI::ShaderStageType::Compute); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1); + commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 2); { commandList->SetPipelineState(calcScalingFactorPipeline_.get()); - commandList->SetConstantBuffer(calcScalingFactorConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(calcScalingFactorConstant_.get(), 0); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); commandList->SetPipelineState(calcCorrectPositionPipeline_.get()); - commandList->SetConstantBuffer(calcCorrectPositionConstant_.get(), LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(calcCorrectPositionConstant_.get(), 0); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); } commandList->ResetComputeBuffer(); } commandList->SetPipelineState(integratePipeline_.get()); - commandList->SetConstantBuffer(integrateConstant_.get(), LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(integrateConstant_.get(), 0); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); commandList->ResetComputeBuffer(); commandList->SetPipelineState(buildVBIBPipeline_.get()); - commandList->SetConstantBuffer(buildVBIBConstant_.get(), LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0, LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(particleComputeVertex_.get(), sizeof(Vertex), 1, LLGI::ShaderStageType::Compute); - commandList->SetComputeBuffer(particleComputeIndex_.get(), sizeof(int), 2, LLGI::ShaderStageType::Compute); + commandList->SetConstantBuffer(buildVBIBConstant_.get(), 0); + commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + commandList->SetComputeBuffer(particleComputeVertex_.get(), sizeof(Vertex), 1); + commandList->SetComputeBuffer(particleComputeIndex_.get(), sizeof(int), 2); commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); commandList->ResetComputeBuffer(); From 7bb29324dbb5c909f96533b2f86de046f16f380c Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 19 Aug 2023 09:01:14 +0900 Subject: [PATCH 09/13] re transpile --- .../Shaders/GLSL_VULKAN/bitonicSort.comp | 2 +- .../Shaders/GLSL_VULKAN/buildGrid.comp | 4 ++-- .../Shaders/GLSL_VULKAN/buildGridIndices.comp | 4 ++-- .../Shaders/GLSL_VULKAN/buildVBIB.comp | 6 +++--- .../GLSL_VULKAN/calcCorrectPosition.comp | 6 +++--- .../GLSL_VULKAN/calcExternalForce.comp | 2 +- .../GLSL_VULKAN/calcScalingFactor.comp | 6 +++--- .../Shaders/GLSL_VULKAN/clearGridIndices.comp | 2 +- .../Shaders/GLSL_VULKAN/integrate.comp | 2 +- .../Fluid2D/Shaders/Metal/bitonicSort.comp | 2 +- examples/Fluid2D/Shaders/Metal/buildGrid.comp | 2 +- .../Shaders/Metal/buildGridIndices.comp | 2 +- examples/Fluid2D/Shaders/Metal/buildVBIB.comp | 2 +- .../Shaders/Metal/calcCorrectPosition.comp | 2 +- .../Shaders/Metal/calcExternalForce.comp | 2 +- .../Shaders/Metal/calcScalingFactor.comp | 2 +- .../Shaders/Metal/clearGridIndices.comp | 2 +- examples/Fluid2D/Shaders/Metal/integrate.comp | 2 +- .../Shaders/SPIRV/bitonicSort.comp.spv | Bin 2700 -> 2700 bytes .../Fluid2D/Shaders/SPIRV/buildGrid.comp.spv | Bin 3532 -> 3532 bytes .../Shaders/SPIRV/buildGridIndices.comp.spv | Bin 2868 -> 2868 bytes .../Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv | Bin 6956 -> 6956 bytes .../SPIRV/calcCorrectPosition.comp.spv | Bin 9456 -> 9456 bytes .../Shaders/SPIRV/calcExternalForce.comp.spv | Bin 3048 -> 3048 bytes .../Shaders/SPIRV/calcScalingFactor.comp.spv | Bin 10812 -> 10812 bytes .../Shaders/SPIRV/clearGridIndices.comp.spv | Bin 1096 -> 1096 bytes .../Fluid2D/Shaders/SPIRV/integrate.comp.spv | Bin 1920 -> 1920 bytes .../Fluid2D/Shaders/SPIRV/render.frag.spv | Bin 2388 -> 2388 bytes .../Fluid2D/Shaders/SPIRV/render.vert.spv | Bin 2776 -> 2776 bytes 29 files changed, 26 insertions(+), 26 deletions(-) diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp index 3eef971..c5ce564 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp @@ -7,7 +7,7 @@ layout(set = 0, binding = 0, std140) uniform CB float Dir; } _19; -layout(set = 0, binding = 1, std430) buffer gridTable +layout(set = 2, binding = 0, std430) buffer gridTable { ivec2 _data[]; } gridTable_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp index 7964108..f2d6fbe 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGrid.comp @@ -17,12 +17,12 @@ layout(set = 0, binding = 0, std140) uniform CB int ParticlesCount; } _52; -layout(set = 0, binding = 2, std430) buffer gridTable +layout(set = 2, binding = 1, std430) buffer gridTable { ivec2 _data[]; } gridTable_1; -layout(set = 0, binding = 1, std430) buffer particles +layout(set = 2, binding = 0, std430) buffer particles { Particle _data[]; } particles_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp index 8920fad..6416368 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildGridIndices.comp @@ -6,12 +6,12 @@ layout(set = 0, binding = 0, std140) uniform CB int ParticlesCount; } _22; -layout(set = 0, binding = 1, std430) buffer gridTable +layout(set = 2, binding = 0, std430) buffer gridTable { ivec2 _data[]; } gridTable_1; -layout(set = 0, binding = 2, std430) buffer gridIndicesTable +layout(set = 2, binding = 1, std430) buffer gridIndicesTable { ivec2 _data[]; } gridIndicesTable_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp index 32d0d47..fe26b15 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp @@ -25,17 +25,17 @@ layout(set = 0, binding = 0, std140) uniform CB vec4 FixedColor; } _67; -layout(set = 0, binding = 1, std430) buffer particles +layout(set = 2, binding = 0, std430) buffer particles { Particle _data[]; } particles_1; -layout(set = 0, binding = 2, std430) buffer vertex_ +layout(set = 2, binding = 1, std430) buffer vertex_ { Vertex _data[]; } vertex_1; -layout(set = 0, binding = 3, std430) buffer index +layout(set = 2, binding = 2, std430) buffer index { int _data[]; } index_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp index 257d97a..a023e3c 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcCorrectPosition.comp @@ -22,17 +22,17 @@ layout(set = 0, binding = 0, std140) uniform CB float GWspiky; } _60; -layout(set = 0, binding = 3, std430) buffer gridIndicesTable +layout(set = 2, binding = 2, std430) buffer gridIndicesTable { ivec2 _data[]; } gridIndicesTable_1; -layout(set = 0, binding = 1, std430) buffer particles +layout(set = 2, binding = 0, std430) buffer particles { Particle _data[]; } particles_1; -layout(set = 0, binding = 2, std430) buffer gridTable +layout(set = 2, binding = 1, std430) buffer gridTable { ivec2 _data[]; } gridTable_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp index e87384b..2acb5fe 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcExternalForce.comp @@ -10,7 +10,7 @@ struct Particle uint IsFix; }; -layout(set = 0, binding = 1, std430) buffer particles +layout(set = 2, binding = 0, std430) buffer particles { Particle _data[]; } particles_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp index 840cfe4..4fc481b 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/calcScalingFactor.comp @@ -22,17 +22,17 @@ layout(set = 0, binding = 0, std140) uniform CB float GWspiky; } _66; -layout(set = 0, binding = 3, std430) buffer gridIndicesTable +layout(set = 2, binding = 2, std430) buffer gridIndicesTable { ivec2 _data[]; } gridIndicesTable_1; -layout(set = 0, binding = 1, std430) buffer particles +layout(set = 2, binding = 0, std430) buffer particles { Particle _data[]; } particles_1; -layout(set = 0, binding = 2, std430) buffer gridTable +layout(set = 2, binding = 1, std430) buffer gridTable { ivec2 _data[]; } gridTable_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp index b6892fb..96f27da 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/clearGridIndices.comp @@ -1,7 +1,7 @@ #version 430 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; -layout(set = 0, binding = 1, std430) buffer gridIndicesTable +layout(set = 2, binding = 0, std430) buffer gridIndicesTable { ivec2 _data[]; } gridIndicesTable_1; diff --git a/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp index abb90bb..ab74f5e 100644 --- a/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp @@ -10,7 +10,7 @@ struct Particle uint IsFix; }; -layout(set = 0, binding = 1, std430) buffer particles +layout(set = 2, binding = 0, std430) buffer particles { Particle _data[]; } particles_1; diff --git a/examples/Fluid2D/Shaders/Metal/bitonicSort.comp b/examples/Fluid2D/Shaders/Metal/bitonicSort.comp index 246fef7..2c7711d 100644 --- a/examples/Fluid2D/Shaders/Metal/bitonicSort.comp +++ b/examples/Fluid2D/Shaders/Metal/bitonicSort.comp @@ -38,7 +38,7 @@ void _main(thread const uint3& dtid, constant CB& _19, device gridTable& gridTab gridTable_1._data[inc + i] = x1; } -kernel void main0(constant CB& _19 [[buffer(0)]], device gridTable& gridTable_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _19 [[buffer(0)]], device gridTable& gridTable_1 [[buffer(10)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/buildGrid.comp b/examples/Fluid2D/Shaders/Metal/buildGrid.comp index 459799f..28b79d9 100644 --- a/examples/Fluid2D/Shaders/Metal/buildGrid.comp +++ b/examples/Fluid2D/Shaders/Metal/buildGrid.comp @@ -62,7 +62,7 @@ void _main(thread const uint3& dtid, constant CB& _52, device gridTable& gridTab } } -kernel void main0(constant CB& _52 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device gridTable& gridTable_1 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _52 [[buffer(0)]], device particles& particles_1 [[buffer(10)]], device gridTable& gridTable_1 [[buffer(11)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp b/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp index 479e1c1..8f29df8 100644 --- a/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp +++ b/examples/Fluid2D/Shaders/Metal/buildGridIndices.comp @@ -49,7 +49,7 @@ void _main(thread const uint3& dtid, constant CB& _22, device gridTable& gridTab } } -kernel void main0(constant CB& _22 [[buffer(0)]], device gridTable& gridTable_1 [[buffer(1)]], device gridTable& gridIndicesTable [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _22 [[buffer(0)]], device gridTable& gridTable_1 [[buffer(10)]], device gridTable& gridIndicesTable [[buffer(11)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/buildVBIB.comp b/examples/Fluid2D/Shaders/Metal/buildVBIB.comp index 1fde583..52a02e7 100644 --- a/examples/Fluid2D/Shaders/Metal/buildVBIB.comp +++ b/examples/Fluid2D/Shaders/Metal/buildVBIB.comp @@ -84,7 +84,7 @@ void _main(thread const uint3& dtid, constant CB& _67, device particles& particl index_1._data[(dtid.x * 6u) + 5u] = int((dtid.x * 4u) + 3u); } -kernel void main0(constant CB& _67 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device vertex_& vertex_1 [[buffer(2)]], device index& index_1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _67 [[buffer(0)]], device particles& particles_1 [[buffer(10)]], device vertex_& vertex_1 [[buffer(11)]], device index& index_1 [[buffer(12)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp b/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp index 1dea653..3efbc2b 100644 --- a/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp +++ b/examples/Fluid2D/Shaders/Metal/calcCorrectPosition.comp @@ -127,7 +127,7 @@ void _main(thread const uint3& dtid, constant CB& _60, device gridIndicesTable& particles_1._data[id].Next += CalcPositionCorrection(param, _60, gridIndicesTable_1, particles_1, gridTable); } -kernel void main0(constant CB& _60 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device gridIndicesTable& gridTable [[buffer(2)]], device gridIndicesTable& gridIndicesTable_1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _60 [[buffer(0)]], device particles& particles_1 [[buffer(10)]], device gridIndicesTable& gridTable [[buffer(11)]], device gridIndicesTable& gridIndicesTable_1 [[buffer(12)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp b/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp index dbda0c9..28b0408 100644 --- a/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp +++ b/examples/Fluid2D/Shaders/Metal/calcExternalForce.comp @@ -42,7 +42,7 @@ void _main(thread const uint3& dtid, device particles& particles_1, constant CB& particles_1._data[dtid.x].Next = pos; } -kernel void main0(constant CB& _41 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _41 [[buffer(0)]], device particles& particles_1 [[buffer(10)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp b/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp index fe82782..3d7f579 100644 --- a/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp +++ b/examples/Fluid2D/Shaders/Metal/calcScalingFactor.comp @@ -147,7 +147,7 @@ void _main(thread const uint3& dtid, constant CB& _66, device gridIndicesTable& CalcScalingFactor(param, _66, gridIndicesTable_1, particles_1, gridTable); } -kernel void main0(constant CB& _66 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], device gridIndicesTable& gridTable [[buffer(2)]], device gridIndicesTable& gridIndicesTable_1 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _66 [[buffer(0)]], device particles& particles_1 [[buffer(10)]], device gridIndicesTable& gridTable [[buffer(11)]], device gridIndicesTable& gridIndicesTable_1 [[buffer(12)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp b/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp index 1fa10db..4d23d2d 100644 --- a/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp +++ b/examples/Fluid2D/Shaders/Metal/clearGridIndices.comp @@ -16,7 +16,7 @@ void _main(thread const uint3& dtid, device gridIndicesTable& gridIndicesTable_1 gridIndicesTable_1._data[dtid.x] = int2(2147483647); } -kernel void main0(device gridIndicesTable& gridIndicesTable_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(device gridIndicesTable& gridIndicesTable_1 [[buffer(10)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/Metal/integrate.comp b/examples/Fluid2D/Shaders/Metal/integrate.comp index 6cfd416..626907d 100644 --- a/examples/Fluid2D/Shaders/Metal/integrate.comp +++ b/examples/Fluid2D/Shaders/Metal/integrate.comp @@ -31,7 +31,7 @@ void _main(thread const uint3& dtid, device particles& particles_1, constant CB& particles_1._data[dtid.x].Current = particles_1._data[dtid.x].Next; } -kernel void main0(constant CB& _40 [[buffer(0)]], device particles& particles_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant CB& _40 [[buffer(0)]], device particles& particles_1 [[buffer(10)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { uint3 dtid = gl_GlobalInvocationID; uint3 param = dtid; diff --git a/examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv b/examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv index 59433053d9a5ef618e42d12a808937973cf618b6..886fe47a21ec34aabcee53e9cdfe3013fd9ff3fb 100644 GIT binary patch delta 41 tcmeAX?GfcRKnU|?k6-pE;P`01;hXV delta 41 scmeAX?GfcRKnU|?k6+Q?bYB**{+?hGsp)<9Yjh#5D}Wcthw0B&Lh!TjeMXv-nMs+Qfq{{Mdn2baiy#vN1A{vQ3xhL|Rs>>($+0ZjeMXv-nMs+Qfq{{MYa^#Kiy#9KxHGUYI0I=#AZDB#%VG}~je&|XZC=YF&j|qj CLQO3=Q%$L~# D2=@o| delta 63 zcmdlYwndDSnMs+Qfq{{MYa{1(CP4-uaA#m)um{qLK+HHNK CV+Zm8 diff --git a/examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv index 520b65e750dd79934f2579ebabc9fade08861179..26bec3d5f29214778760dff70ad41f234526817e 100644 GIT binary patch delta 85 zcmZ2uw#JN;nMs+Qfq{{Mdn0E&iy#vN1A{vQ3qt^qRs>>($(1bqaM1#&DC1;FR(rVU NRj4S_=1kVZq5#UN3c~;Z delta 85 zcmZ2uw#JN;nMs+Qfq{{MYa?eoiy#9KxHGUY1ORD8AZDCg$?E$+nz};iA=0QHIGMIql)1 N^-xj9&5m5Z>(%`+L-GXVg4hy|_y delta 41 scmX@Xae{-BnMs+Qfq{{MYa?epqaXtixHGUY2mxtDAZFY=lW{#00D49RtN;K2 diff --git a/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv b/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv index 5d2a966a54d6e7115cd06e779efbfe4b75c77012..bc42d7864dafef3c77f33a9afec50c1ae960984b 100644 GIT binary patch delta 41 tcmZqRZ{X)-W>RKnU|?k6-pERKnU|?k6+Q?bXB**{+?hGsp!a!OPh#5CeWO~m60A;HMUjP6A diff --git a/examples/Fluid2D/Shaders/SPIRV/render.frag.spv b/examples/Fluid2D/Shaders/SPIRV/render.frag.spv index 275e7ec4d2af9bf99172dcc0553e7795a65b5227..f7db3c1bd4310d1babfa502c78918c469220cc62 100644 GIT binary patch delta 18 Zcmca2bVZ1hnMs+Qfq{{Mdn2beCjcb610MhY delta 18 Zcmca2bVZ1hnMs+Qfq{{MYa^#OCjcb110DbX diff --git a/examples/Fluid2D/Shaders/SPIRV/render.vert.spv b/examples/Fluid2D/Shaders/SPIRV/render.vert.spv index fe8e5d88b165633c37ff3492c7dcb374dde1cc4e..804af046a329d58c9c6f850ead0e0b1cb120df29 100644 GIT binary patch delta 18 Zcmca1dP9_xnMs+Qfq{{Mdn4z0E&wL01F`@B delta 18 Zcmca1dP9_xnMs+Qfq{{MYa{1*E&wK`1F--A From 6ee932fcd5faa1e5b9824ade34b04ce81e27a3ee Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 19 Aug 2023 11:02:31 +0900 Subject: [PATCH 10/13] add assertion --- src/DX12/LLGI.CommandListDX12.cpp | 1 + 1 file changed, 1 insertion(+) diff --git a/src/DX12/LLGI.CommandListDX12.cpp b/src/DX12/LLGI.CommandListDX12.cpp index 2e8ed73..9f04309 100644 --- a/src/DX12/LLGI.CommandListDX12.cpp +++ b/src/DX12/LLGI.CommandListDX12.cpp @@ -308,6 +308,7 @@ void CommandListDX12::Draw(int32_t primitiveCount, int32_t instanceCount) { currentCommandList_->SetGraphicsRootSignature(pip->GetRootSignature()); auto p = pip->GetPipelineState(); + assert(p != nullptr); currentCommandList_->SetPipelineState(p); currentCommandList_->OMSetStencilRef(pip->StencilRef); } From 8e647cb0dd60bc8ccd97dd7f1b1d6148151863ed Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 19 Aug 2023 11:03:19 +0900 Subject: [PATCH 11/13] fix render --- examples/Fluid2D/Fluid2D.cpp | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/examples/Fluid2D/Fluid2D.cpp b/examples/Fluid2D/Fluid2D.cpp index 08971f2..beff0e9 100644 --- a/examples/Fluid2D/Fluid2D.cpp +++ b/examples/Fluid2D/Fluid2D.cpp @@ -559,14 +559,18 @@ void Fluid2D::Render(LLGI::Graphics* graphics, LLGI::CommandList* commandList, L if (itr == pipelineCache_.end()) { const auto pip = graphics->CreatePiplineState(); - pip->VertexLayouts[0] = LLGI::VertexLayoutFormat::R32G32B32_FLOAT; + pip->VertexLayouts[0] = LLGI::VertexLayoutFormat::R32G32B32A32_FLOAT; pip->VertexLayouts[1] = LLGI::VertexLayoutFormat::R8G8B8A8_UNORM; + pip->VertexLayouts[2] = LLGI::VertexLayoutFormat::R32G32_FLOAT; pip->VertexLayouts[3] = LLGI::VertexLayoutFormat::R32G32_FLOAT; - pip->VertexLayouts[4] = LLGI::VertexLayoutFormat::R32G32_FLOAT; pip->VertexLayoutNames[0] = "POSITION"; pip->VertexLayoutNames[1] = "COLOR"; - pip->VertexLayoutNames[2] = "UV1"; - pip->VertexLayoutNames[3] = "UV2"; + pip->VertexLayoutNames[2] = "UV"; + pip->VertexLayoutNames[3] = "UV"; + pip->VertexLayoutSemantics[0] = 0; + pip->VertexLayoutSemantics[1] = 0; + pip->VertexLayoutSemantics[2] = 0; + pip->VertexLayoutSemantics[3] = 1; pip->VertexLayoutCount = 4; pip->Topology = LLGI::TopologyType::Triangle; @@ -581,7 +585,7 @@ void Fluid2D::Render(LLGI::Graphics* graphics, LLGI::CommandList* commandList, L pip->SetShader(LLGI::ShaderStageType::Vertex, vs_.get()); pip->SetShader(LLGI::ShaderStageType::Pixel, ps_.get()); pip->SetRenderPassPipelineState(renderPassState.get()); - pip->Compile(); + assert(pip->Compile()); pipeline = LLGI::CreateSharedPtr(pip); pipelineCache_[renderPassState] = pipeline; } From 561ef98dcc074b2c8edb49f3274ae632c72f3ce3 Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 19 Aug 2023 11:03:36 +0900 Subject: [PATCH 12/13] update with assert --- examples/Fluid2D/Fluid2D.cpp | 66 +++++++++++------------------------- 1 file changed, 19 insertions(+), 47 deletions(-) diff --git a/examples/Fluid2D/Fluid2D.cpp b/examples/Fluid2D/Fluid2D.cpp index beff0e9..dcf7e99 100644 --- a/examples/Fluid2D/Fluid2D.cpp +++ b/examples/Fluid2D/Fluid2D.cpp @@ -63,7 +63,11 @@ std::shared_ptr CreateShader( data.push_back(d); - return LLGI::CreateSharedPtr(graphics->CreateShader(data.data(), static_cast(data.size()))); + const auto shader = graphics->CreateShader(data.data(), static_cast(data.size())); + + assert(shader != nullptr); + + return LLGI::CreateSharedPtr(shader); } else { @@ -85,7 +89,11 @@ std::shared_ptr CreateShader( data.push_back(d); } - return LLGI::CreateSharedPtr(graphics->CreateShader(data.data(), static_cast(data.size()))); + const auto shader = graphics->CreateShader(data.data(), static_cast(data.size())); + + assert(shader != nullptr); + + return LLGI::CreateSharedPtr(shader); } } @@ -191,11 +199,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(CalcExternalForceCB)) ); calcExternalPipeline_->SetShader(LLGI::ShaderStageType::Compute, calcExternalShader_.get()); - if (!calcExternalPipeline_->Compile()) - { - std::cerr << "failed to compile calcExternalPipline" << std::endl; - abort(); - } + assert(calcExternalPipeline_->Compile()); { const auto cb = (CalcExternalForceCB*)calcExternalConstant_->Lock(); @@ -220,11 +224,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) ); buildGridPipeline_->SetShader(LLGI::ShaderStageType::Compute, buildGridShader_.get()); - if (!buildGridPipeline_->Compile()) - { - std::cerr << "failed to compile buildGridPipeline" << std::endl; - abort(); - } + assert(buildGridPipeline_->Compile()); { const auto cb = (BuildGridCB*)buildGridConstant_->Lock(); @@ -244,21 +244,13 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) ); bitonicSortPipeline_->SetShader(LLGI::ShaderStageType::Compute, bitonicSortShader_.get()); - if (!bitonicSortPipeline_->Compile()) - { - std::cerr << "failed to compile bitonicSortPipeline" << std::endl; - abort(); - } + assert(bitonicSortPipeline_->Compile()); // clearGridIndices std::cout << "Creating clearGridIndicesPipeline_" << std::endl; clearGridIndicesPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); clearGridIndicesPipeline_->SetShader(LLGI::ShaderStageType::Compute, clearGridIndicesShader_.get()); - if (!clearGridIndicesPipeline_->Compile()) - { - std::cerr << "failed to compile clearGridIndicesPipeline" << std::endl; - abort(); - } + assert(clearGridIndicesPipeline_->Compile()); // buildGridIndices std::cout << "Creating buildGridIndicesPipeline_" << std::endl; @@ -267,11 +259,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(BuildGridIndicesCB)) ); buildGridIndicesPipeline_->SetShader(LLGI::ShaderStageType::Compute, buildGridIndicesShader_.get()); - if (!buildGridIndicesPipeline_->Compile()) - { - std::cerr << "failed to compile buildGridIndicesPipeline" << std::endl; - abort(); - } + assert(buildGridIndicesPipeline_->Compile()); { const auto cb = (BuildGridIndicesCB*)buildGridIndicesConstant_->Lock(); @@ -288,11 +276,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) graphics->CreateBuffer(LLGI::BufferUsageType::Constant | LLGI::BufferUsageType::MapWrite, sizeof(CalcScalingFactorCB)) ); calcScalingFactorPipeline_->SetShader(LLGI::ShaderStageType::Compute, calcScalingFactorShader_.get()); - if (!calcScalingFactorPipeline_->Compile()) - { - std::cerr << "failed to compile calcScalingFactorPipeline" << std::endl; - abort(); - } + assert(calcScalingFactorPipeline_->Compile()); { const auto cb = (CalcScalingFactorCB*)calcScalingFactorConstant_->Lock(); @@ -317,11 +301,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) ); calcCorrectPositionPipeline_->SetShader(LLGI::ShaderStageType::Compute, calcCorrectPositionShader_.get()); - if (!calcCorrectPositionPipeline_->Compile()) - { - std::cerr << "failed to compile calcCorrectPositionPipeline" << std::endl; - abort(); - } + assert(calcCorrectPositionPipeline_->Compile()); { const auto cb = (CalcCorrectPositionCB*)calcCorrectPositionConstant_->Lock(); @@ -346,11 +326,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) ); integratePipeline_->SetShader(LLGI::ShaderStageType::Compute, integrateShader_.get()); - if (!integratePipeline_->Compile()) - { - std::cerr << "failed to compile integratePipeline" << std::endl; - abort(); - } + assert(integratePipeline_->Compile()); { const auto cb = (IntegrateCB*)integrateConstant_->Lock(); @@ -368,11 +344,7 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) ); buildVBIBPipeline_->SetShader(LLGI::ShaderStageType::Compute, buildVBIBShader_.get()); - if (!buildVBIBPipeline_->Compile()) - { - std::cerr << "failed to compile buildVBIBPipeline" << std::endl; - abort(); - } + assert(buildVBIBPipeline_->Compile()); { const auto cb = (BuildVBIBCB*)buildVBIBConstant_->Lock(); From 047db7c72b253f6d60460c1e302a21b1fbf20c75 Mon Sep 17 00:00:00 2001 From: wraikny Date: Sat, 19 Aug 2023 12:16:31 +0900 Subject: [PATCH 13/13] update --- examples/Fluid2D/Fluid2D.cpp | 168 +++++++++--------- examples/Fluid2D/Fluid2D.h | 2 +- .../Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp | 26 +-- .../Fluid2D/Shaders/HLSL_DX12/render.frag | 8 +- .../Fluid2D/Shaders/HLSL_DX12/render.vert | 11 +- 5 files changed, 102 insertions(+), 113 deletions(-) diff --git a/examples/Fluid2D/Fluid2D.cpp b/examples/Fluid2D/Fluid2D.cpp index dcf7e99..5964754 100644 --- a/examples/Fluid2D/Fluid2D.cpp +++ b/examples/Fluid2D/Fluid2D.cpp @@ -350,15 +350,15 @@ Fluid2D::Fluid2D(LLGI::Graphics* graphics, LLGI::DeviceType deviceType) const auto cb = (BuildVBIBCB*)buildVBIBConstant_->Lock(); cb->ParticleRadius = ParticleRadius; - cb->Color[0] = 0.2f; - cb->Color[1] = 0.2f; - cb->Color[2] = 1.0f; - cb->Color[3] = 1.0f; + cb->Color[0] = 50u; + cb->Color[1] = 50u; + cb->Color[2] = 255u; + cb->Color[3] = 255u; - cb->FixedColor[0] = 0.0f; - cb->FixedColor[1] = 0.0f; - cb->FixedColor[2] = 0.0f; - cb->FixedColor[3] = 1.0f; + cb->FixedColor[0] = 0u; + cb->FixedColor[1] = 0u; + cb->FixedColor[2] = 0u; + cb->FixedColor[3] = 255u; buildVBIBConstant_->Unlock(); } @@ -428,81 +428,81 @@ void Fluid2D::Update(LLGI::Graphics* graphics, LLGI::CommandList* commandList) commandList->BeginComputePass(); - commandList->SetPipelineState(calcExternalPipeline_.get()); - commandList->SetConstantBuffer(calcExternalConstant_.get(), 0); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); - commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); - commandList->ResetComputeBuffer(); - - for (int l = 0; l < IterationCount; l++) - { - commandList->SetPipelineState(buildGridPipeline_.get()); - commandList->SetConstantBuffer(buildGridConstant_.get(), 0); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1); - commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); - commandList->ResetComputeBuffer(); - - const int nlog = static_cast(std::log2f(ParticlesCount)); - int inc; - - commandList->SetPipelineState(bitonicSortPipeline_.get()); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0); - for (int i = 0; i < nlog; i++) - { - inc = 1 << i; - - const auto dir = 2 << i; - - for (int j = 0; j < i + 1; j++) - { - { - const auto cb = (BitonicSortCB*)bitonicSortConstant_->Lock(); - cb->Inc = inc; - cb->Dir = dir; - bitonicSortConstant_->Unlock(); - } - - commandList->SetConstantBuffer(bitonicSortConstant_.get(), 0); - commandList->Dispatch(ParticlesCount / 2, 1, 1, 1, 1, 1); - - inc /= 2; - } - } - commandList->ResetComputeBuffer(); - - commandList->SetPipelineState(clearGridIndicesPipeline_.get()); - commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 0); - commandList->Dispatch(gridNum_.X * gridNum_.Y, 1, 1, 1, 1, 1); - commandList->ResetComputeBuffer(); - - commandList->SetPipelineState(buildGridIndicesPipeline_.get()); - commandList->SetConstantBuffer(buildGridIndicesConstant_.get(), 0); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0); - commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 1); - commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); - commandList->ResetComputeBuffer(); - - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); - commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1); - commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 2); - { - commandList->SetPipelineState(calcScalingFactorPipeline_.get()); - commandList->SetConstantBuffer(calcScalingFactorConstant_.get(), 0); - commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); - - commandList->SetPipelineState(calcCorrectPositionPipeline_.get()); - commandList->SetConstantBuffer(calcCorrectPositionConstant_.get(), 0); - commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); - } - commandList->ResetComputeBuffer(); - } - - commandList->SetPipelineState(integratePipeline_.get()); - commandList->SetConstantBuffer(integrateConstant_.get(), 0); - commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); - commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); - commandList->ResetComputeBuffer(); + // commandList->SetPipelineState(calcExternalPipeline_.get()); + // commandList->SetConstantBuffer(calcExternalConstant_.get(), 0); + // commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + // commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + // commandList->ResetComputeBuffer(); + + // for (int l = 0; l < IterationCount; l++) + // { + // commandList->SetPipelineState(buildGridPipeline_.get()); + // commandList->SetConstantBuffer(buildGridConstant_.get(), 0); + // commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + // commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1); + // commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + // commandList->ResetComputeBuffer(); + + // const int nlog = static_cast(std::log2f(ParticlesCount)); + // int inc; + + // commandList->SetPipelineState(bitonicSortPipeline_.get()); + // commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0); + // for (int i = 0; i < nlog; i++) + // { + // inc = 1 << i; + + // const auto dir = 2 << i; + + // for (int j = 0; j < i + 1; j++) + // { + // { + // const auto cb = (BitonicSortCB*)bitonicSortConstant_->Lock(); + // cb->Inc = inc; + // cb->Dir = dir; + // bitonicSortConstant_->Unlock(); + // } + + // commandList->SetConstantBuffer(bitonicSortConstant_.get(), 0); + // commandList->Dispatch(ParticlesCount / 2, 1, 1, 1, 1, 1); + + // inc /= 2; + // } + // } + // commandList->ResetComputeBuffer(); + + // commandList->SetPipelineState(clearGridIndicesPipeline_.get()); + // commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 0); + // commandList->Dispatch(gridNum_.X * gridNum_.Y, 1, 1, 1, 1, 1); + // commandList->ResetComputeBuffer(); + + // commandList->SetPipelineState(buildGridIndicesPipeline_.get()); + // commandList->SetConstantBuffer(buildGridIndicesConstant_.get(), 0); + // commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 0); + // commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 1); + // commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + // commandList->ResetComputeBuffer(); + + // commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + // commandList->SetComputeBuffer(gridTable_.get(), sizeof(int[2]), 1); + // commandList->SetComputeBuffer(gridIndicesTable_.get(), sizeof(int[2]), 2); + // { + // commandList->SetPipelineState(calcScalingFactorPipeline_.get()); + // commandList->SetConstantBuffer(calcScalingFactorConstant_.get(), 0); + // commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + + // commandList->SetPipelineState(calcCorrectPositionPipeline_.get()); + // commandList->SetConstantBuffer(calcCorrectPositionConstant_.get(), 0); + // commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + // } + // commandList->ResetComputeBuffer(); + // } + + // commandList->SetPipelineState(integratePipeline_.get()); + // commandList->SetConstantBuffer(integrateConstant_.get(), 0); + // commandList->SetComputeBuffer(particles_.get(), sizeof(Particle), 0); + // commandList->Dispatch(ParticlesCount, 1, 1, 1, 1, 1); + // commandList->ResetComputeBuffer(); commandList->SetPipelineState(buildVBIBPipeline_.get()); commandList->SetConstantBuffer(buildVBIBConstant_.get(), 0); @@ -531,7 +531,7 @@ void Fluid2D::Render(LLGI::Graphics* graphics, LLGI::CommandList* commandList, L if (itr == pipelineCache_.end()) { const auto pip = graphics->CreatePiplineState(); - pip->VertexLayouts[0] = LLGI::VertexLayoutFormat::R32G32B32A32_FLOAT; + pip->VertexLayouts[0] = LLGI::VertexLayoutFormat::R32G32B32_FLOAT; pip->VertexLayouts[1] = LLGI::VertexLayoutFormat::R8G8B8A8_UNORM; pip->VertexLayouts[2] = LLGI::VertexLayoutFormat::R32G32_FLOAT; pip->VertexLayouts[3] = LLGI::VertexLayoutFormat::R32G32_FLOAT; diff --git a/examples/Fluid2D/Fluid2D.h b/examples/Fluid2D/Fluid2D.h index 9487d65..d301fc3 100644 --- a/examples/Fluid2D/Fluid2D.h +++ b/examples/Fluid2D/Fluid2D.h @@ -62,7 +62,7 @@ struct alignas(16) BuildVBIBCB struct alignas(16) Vertex { float Position[3]; - float Color[3]; + float Color[4]; float UV1[2]; float UV2[2]; }; diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp index 578efe8..859f13e 100644 --- a/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp +++ b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp @@ -22,7 +22,7 @@ int2 GetGridPos(float2 pos, float2 gridSize) struct Vertex { float3 Position; - int Color; + float4 Color; float2 UV1; float2 UV2; }; @@ -38,31 +38,19 @@ RWStructuredBuffer particles : register(u0); RWStructuredBuffer vertex_ : register(u1); RWStructuredBuffer index : register(u2); -int DecodeFloatRGBA(float4 rgba) { - int res = 0; - res += (int)(rgba.a * 255); - res = res << 8; - res += (int)(rgba.b * 255); - res = res << 8; - res += (int)(rgba.g * 255); - res = res << 8; - res += (int)(rgba.r * 255); - return res; -} - [numthreads(1, 1, 1)] void main(uint3 dtid : SV_DispatchThreadID) { - int c = DecodeFloatRGBA(lerp(Color, FixedColor, particles[dtid.x].IsFix)); + float4 c = lerp(Color, FixedColor, particles[dtid.x].IsFix); for (int i = 0; i < 4; i++) { vertex_[dtid.x * 4 + i].Color = c; } - float3 pos = float3(particles[dtid.x].Current, 0.5); - vertex_[dtid.x * 4].Position = pos + float3(-1, -1, 0) * ParticleRadius * 3; - vertex_[dtid.x * 4 + 1].Position = pos + float3(1, -1, 0) * ParticleRadius * 3; - vertex_[dtid.x * 4 + 2].Position = pos + float3(1, 1, 0) * ParticleRadius * 3; - vertex_[dtid.x * 4 + 3].Position = pos + float3(-1, 1, 0) * ParticleRadius * 3; + float3 pos = float3(particles[dtid.x].Current, 0.0); + vertex_[dtid.x * 4].Position = pos + float3(-1, -1, 0) * ParticleRadius; + vertex_[dtid.x * 4 + 1].Position = pos + float3(1, -1, 0) * ParticleRadius; + vertex_[dtid.x * 4 + 2].Position = pos + float3(1, 1, 0) * ParticleRadius; + vertex_[dtid.x * 4 + 3].Position = pos + float3(-1, 1, 0) * ParticleRadius; vertex_[dtid.x * 4].UV1 = float2(0, 0); vertex_[dtid.x * 4 + 1].UV1 = float2(1, 0); vertex_[dtid.x * 4 + 2].UV1 = float2(1, 1); diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/render.frag b/examples/Fluid2D/Shaders/HLSL_DX12/render.frag index 9e41453..5762e4e 100644 --- a/examples/Fluid2D/Shaders/HLSL_DX12/render.frag +++ b/examples/Fluid2D/Shaders/HLSL_DX12/render.frag @@ -1,9 +1,9 @@ struct PS_INPUT { - float4 Position : SV_POSITION; - float4 Color : COLOR0; - float2 UV1 : UV0; - float2 UV2 : UV1; + float4 Position : SV_POSITION; + float4 Color : COLOR0; + float2 UV1 : UV0; + float2 UV2 : UV1; }; float4 main(PS_INPUT input) : SV_TARGET diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/render.vert b/examples/Fluid2D/Shaders/HLSL_DX12/render.vert index 4b01499..2ea7eeb 100644 --- a/examples/Fluid2D/Shaders/HLSL_DX12/render.vert +++ b/examples/Fluid2D/Shaders/HLSL_DX12/render.vert @@ -6,15 +6,16 @@ struct VS_INPUT{ }; struct VS_OUTPUT{ - float4 Position : SV_POSITION; - float4 Color : COLOR0; - float2 UV1 : UV0; - float2 UV2 : UV1; + float4 Position : SV_POSITION; + float4 Color : COLOR0; + float2 UV1 : UV0; + float2 UV2 : UV1; }; VS_OUTPUT main(VS_INPUT input){ VS_OUTPUT output; - output.Position = float4(input.Position, 1.0f); + output.Position = float4(input.Position / 200.0f - float3(0.5f, 0.5f, 0), 1.0f); + // output.Position = float4(input.Position, 1.0f); output.UV1 = input.UV1; output.UV2 = input.UV2; output.Color = input.Color;