diff --git a/CMakeLists.txt b/CMakeLists.txt index 26a3aba..6ef8237 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -79,7 +79,6 @@ function(add_shader FILE) package(${SPV}) endfunction() add_shader(add1.comp) -add_shader(add2.comp src/config.hpp) add_shader(advect1.comp src/config.hpp shaders/shader.glsl) add_shader(advect2.comp src/config.hpp shaders/shader.glsl) add_shader(bnd1.comp src/config.hpp) diff --git a/samples/1.json b/samples/1.json index 8f5daaf..9958436 100644 --- a/samples/1.json +++ b/samples/1.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.009999999776482582, + "Diffusion": 0.000004, "Iterations": 5, "Size": 128, "Spawners": [ @@ -22,5 +22,5 @@ "Value": 1.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/10.json b/samples/10.json index 427e372..6046720 100644 --- a/samples/10.json +++ b/samples/10.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.009999999776482582, + "Diffusion": 0.000004, "Iterations": 5, "Size": 128, "Spawners": [ @@ -31,5 +31,5 @@ "Value": 8.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/11.json b/samples/11.json index e08b551..7306337 100644 --- a/samples/11.json +++ b/samples/11.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 5, "Size": 128, "Spawners": [ @@ -31,5 +31,5 @@ "Value": -618.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/2.json b/samples/2.json index 1e6023f..32223ca 100644 --- a/samples/2.json +++ b/samples/2.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 10, "Size": 128, "Spawners": [ @@ -31,5 +31,5 @@ "Value": 2.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/3.json b/samples/3.json index ddc89a1..34a5f7f 100644 --- a/samples/3.json +++ b/samples/3.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 10, "Size": 128, "Spawners": [ @@ -31,5 +31,5 @@ "Value": 1.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/4.json b/samples/4.json index 4bfa15c..a8c47d7 100644 --- a/samples/4.json +++ b/samples/4.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 7, "Size": 128, "Spawners": [ @@ -31,5 +31,5 @@ "Value": 2.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/5.json b/samples/5.json index d6facd4..b46f3c1 100644 --- a/samples/5.json +++ b/samples/5.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 5, "Size": 128, "Spawners": [ @@ -40,5 +40,5 @@ "Value": 65.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/6.json b/samples/6.json index 050be28..6dbeb53 100644 --- a/samples/6.json +++ b/samples/6.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 7, "Size": 128, "Spawners": [ @@ -40,5 +40,5 @@ "Value": 1.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/7.json b/samples/7.json index 8708c9f..1148294 100644 --- a/samples/7.json +++ b/samples/7.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.08799999952316284, + "Diffusion": 0.0000352, "Iterations": 3, "Size": 128, "Spawners": [ @@ -40,5 +40,5 @@ "Value": 2.0 } ], - "Viscosity": 0.16500000655651093 + "Viscosity": 0.000066 } \ No newline at end of file diff --git a/samples/8.json b/samples/8.json index b3bff69..e960efc 100644 --- a/samples/8.json +++ b/samples/8.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 7, "Size": 128, "Spawners": [ @@ -40,5 +40,5 @@ "Value": -2.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/samples/9.json b/samples/9.json index ad8e139..fb5261e 100644 --- a/samples/9.json +++ b/samples/9.json @@ -1,5 +1,5 @@ { - "Diffusion": 0.12800000607967377, + "Diffusion": 0.0000512, "Iterations": 5, "Size": 128, "Spawners": [ @@ -49,5 +49,5 @@ "Value": 2.0 } ], - "Viscosity": 0.009999999776482582 + "Viscosity": 0.000004 } \ No newline at end of file diff --git a/shaders/add1.comp b/shaders/add1.comp index 0bbe17f..e29ea8c 100644 --- a/shaders/add1.comp +++ b/shaders/add1.comp @@ -1,7 +1,7 @@ #version 450 layout(local_size_x = 1) in; -layout(set = 1, binding = 0, rgba32f) uniform image3D inOutImage; +layout(set = 1, binding = 0, r32f) uniform image3D inOutImage; layout(set = 2, binding = 0) uniform uniformPosition { ivec3 position; diff --git a/shaders/add2.comp b/shaders/add2.comp deleted file mode 100644 index 666b21e..0000000 --- a/shaders/add2.comp +++ /dev/null @@ -1,22 +0,0 @@ -#version 450 - -#include "config.hpp" - -layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) in; -layout(set = 1, binding = 0, rgba32f) uniform image3D inOutImage; -layout(set = 2, binding = 0) uniform uniformValue -{ - float value; -}; - -void main() -{ - ivec3 size = imageSize(inOutImage); - ivec3 id = ivec3(gl_GlobalInvocationID); - if (any(greaterThanEqual(id, size - 1)) || any(lessThanEqual(id, ivec3(0)))) - { - return; - } - float current = imageLoad(inOutImage, id).x; - imageStore(inOutImage, id, vec4(current + value)); -} \ No newline at end of file diff --git a/shaders/advect1.comp b/shaders/advect1.comp index 051efab..e6ea88d 100644 --- a/shaders/advect1.comp +++ b/shaders/advect1.comp @@ -7,7 +7,7 @@ layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) i layout(set = 0, binding = 0) uniform sampler3D inVelocityX; layout(set = 0, binding = 1) uniform sampler3D inVelocityY; layout(set = 0, binding = 2) uniform sampler3D inVelocityZ; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outVelocity; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outVelocity; layout(set = 2, binding = 0) uniform uniformVelocity { uint velocity; diff --git a/shaders/advect2.comp b/shaders/advect2.comp index b95de5c..ebdb66a 100644 --- a/shaders/advect2.comp +++ b/shaders/advect2.comp @@ -8,7 +8,7 @@ layout(set = 0, binding = 0) uniform sampler3D inDensity; layout(set = 0, binding = 1) uniform sampler3D inVelocityX; layout(set = 0, binding = 2) uniform sampler3D inVelocityY; layout(set = 0, binding = 3) uniform sampler3D inVelocityZ; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outDensity; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outDensity; layout(set = 2, binding = 0) uniform uniformDeltaTime { float deltaTime; diff --git a/shaders/bin/add1.comp.spv b/shaders/bin/add1.comp.spv index 5373192..c1a590e 100644 Binary files a/shaders/bin/add1.comp.spv and b/shaders/bin/add1.comp.spv differ diff --git a/shaders/bin/add2.comp.json b/shaders/bin/add2.comp.json deleted file mode 100644 index a611bcf..0000000 --- a/shaders/bin/add2.comp.json +++ /dev/null @@ -1 +0,0 @@ -{ "samplers": 0, "readonly_storage_textures": 0, "readonly_storage_buffers": 0, "readwrite_storage_textures": 1, "readwrite_storage_buffers": 0, "uniform_buffers": 1, "threadcount_x": 8, "threadcount_y": 8, "threadcount_z": 8 } diff --git a/shaders/bin/add2.comp.msl b/shaders/bin/add2.comp.msl deleted file mode 100644 index cc44e04..0000000 --- a/shaders/bin/add2.comp.msl +++ /dev/null @@ -1,40 +0,0 @@ -#pragma clang diagnostic ignored "-Wmissing-prototypes" - -#include -#include - -using namespace metal; - -template -void spvImageFence(ImageT img) { img.fence(); } - -struct uniformValue -{ - float value; -}; - -constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u); - -kernel void main0(constant uniformValue& _57 [[buffer(0)]], texture3d inOutImage [[texture(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) -{ - int3 size = int3(inOutImage.get_width(), inOutImage.get_height(), inOutImage.get_depth()); - int3 id = int3(gl_GlobalInvocationID); - bool _31 = any(id >= (size - int3(1))); - bool _40; - if (!_31) - { - _40 = any(id <= int3(0)); - } - else - { - _40 = _31; - } - if (_40) - { - return; - } - spvImageFence(inOutImage); - float current = inOutImage.read(uint3(id)).x; - inOutImage.write(float4(current + _57.value), uint3(id)); -} - diff --git a/shaders/bin/add2.comp.spv b/shaders/bin/add2.comp.spv deleted file mode 100644 index ae34478..0000000 Binary files a/shaders/bin/add2.comp.spv and /dev/null differ diff --git a/shaders/bin/advect1.comp.msl b/shaders/bin/advect1.comp.msl index 6877daf..4842204 100644 --- a/shaders/bin/advect1.comp.msl +++ b/shaders/bin/advect1.comp.msl @@ -38,9 +38,9 @@ kernel void main0(constant uniformVelocity& _48 [[buffer(0)]], constant uniformD do { float N = float(size.x - 2); - float dtx = _71.deltaTime * (N - 2.0); - float dty = _71.deltaTime * (N - 2.0); - float dtz = _71.deltaTime * (N - 2.0); + float dtx = _71.deltaTime * N; + float dty = _71.deltaTime * N; + float dtz = _71.deltaTime * N; float tmp1 = dtx * inVelocityX.read(uint3(id), 0).x; float tmp2 = dty * inVelocityY.read(uint3(id), 0).x; float tmp3 = dtz * inVelocityZ.read(uint3(id), 0).x; @@ -65,8 +65,8 @@ kernel void main0(constant uniformVelocity& _48 [[buffer(0)]], constant uniformD int j1i = int(j1); int k0i = int(k0); int k1i = int(k1); - float4 _310 = float4((s0 * ((t0 * ((u0 * inVelocityX.read(uint3(int3(i0i, j0i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i0i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inVelocityX.read(uint3(int3(i0i, j1i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i0i, j1i, k1i)), 0).x))))) + (s1 * ((t0 * ((u0 * inVelocityX.read(uint3(int3(i1i, j0i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i1i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inVelocityX.read(uint3(int3(i1i, j1i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i1i, j1i, k1i)), 0).x)))))); - outVelocity.write(_310, uint3(id)); + float4 _306 = float4((s0 * ((t0 * ((u0 * inVelocityX.read(uint3(int3(i0i, j0i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i0i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inVelocityX.read(uint3(int3(i0i, j1i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i0i, j1i, k1i)), 0).x))))) + (s1 * ((t0 * ((u0 * inVelocityX.read(uint3(int3(i1i, j0i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i1i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inVelocityX.read(uint3(int3(i1i, j1i, k0i)), 0).x) + (u1 * inVelocityX.read(uint3(int3(i1i, j1i, k1i)), 0).x)))))); + outVelocity.write(_306, uint3(id)); } while (false); } else @@ -76,9 +76,9 @@ kernel void main0(constant uniformVelocity& _48 [[buffer(0)]], constant uniformD do { float N_1 = float(size.x - 2); - float dtx_1 = _71.deltaTime * (N_1 - 2.0); - float dty_1 = _71.deltaTime * (N_1 - 2.0); - float dtz_1 = _71.deltaTime * (N_1 - 2.0); + float dtx_1 = _71.deltaTime * N_1; + float dty_1 = _71.deltaTime * N_1; + float dtz_1 = _71.deltaTime * N_1; float tmp1_1 = dtx_1 * inVelocityX.read(uint3(id), 0).x; float tmp2_1 = dty_1 * inVelocityY.read(uint3(id), 0).x; float tmp3_1 = dtz_1 * inVelocityZ.read(uint3(id), 0).x; @@ -103,8 +103,8 @@ kernel void main0(constant uniformVelocity& _48 [[buffer(0)]], constant uniformD int j1i_1 = int(j1_1); int k0i_1 = int(k0_1); int k1i_1 = int(k1_1); - float4 _554 = float4((s0_1 * ((t0_1 * ((u0_1 * inVelocityY.read(uint3(int3(i0i_1, j0i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i0i_1, j0i_1, k1i_1)), 0).x))) + (t1_1 * ((u0_1 * inVelocityY.read(uint3(int3(i0i_1, j1i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i0i_1, j1i_1, k1i_1)), 0).x))))) + (s1_1 * ((t0_1 * ((u0_1 * inVelocityY.read(uint3(int3(i1i_1, j0i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i1i_1, j0i_1, k1i_1)), 0).x))) + (t1_1 * ((u0_1 * inVelocityY.read(uint3(int3(i1i_1, j1i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i1i_1, j1i_1, k1i_1)), 0).x)))))); - outVelocity.write(_554, uint3(id)); + float4 _547 = float4((s0_1 * ((t0_1 * ((u0_1 * inVelocityY.read(uint3(int3(i0i_1, j0i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i0i_1, j0i_1, k1i_1)), 0).x))) + (t1_1 * ((u0_1 * inVelocityY.read(uint3(int3(i0i_1, j1i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i0i_1, j1i_1, k1i_1)), 0).x))))) + (s1_1 * ((t0_1 * ((u0_1 * inVelocityY.read(uint3(int3(i1i_1, j0i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i1i_1, j0i_1, k1i_1)), 0).x))) + (t1_1 * ((u0_1 * inVelocityY.read(uint3(int3(i1i_1, j1i_1, k0i_1)), 0).x) + (u1_1 * inVelocityY.read(uint3(int3(i1i_1, j1i_1, k1i_1)), 0).x)))))); + outVelocity.write(_547, uint3(id)); } while (false); } else @@ -114,9 +114,9 @@ kernel void main0(constant uniformVelocity& _48 [[buffer(0)]], constant uniformD do { float N_2 = float(size.x - 2); - float dtx_2 = _71.deltaTime * (N_2 - 2.0); - float dty_2 = _71.deltaTime * (N_2 - 2.0); - float dtz_2 = _71.deltaTime * (N_2 - 2.0); + float dtx_2 = _71.deltaTime * N_2; + float dty_2 = _71.deltaTime * N_2; + float dtz_2 = _71.deltaTime * N_2; float tmp1_2 = dtx_2 * inVelocityX.read(uint3(id), 0).x; float tmp2_2 = dty_2 * inVelocityY.read(uint3(id), 0).x; float tmp3_2 = dtz_2 * inVelocityZ.read(uint3(id), 0).x; @@ -141,8 +141,8 @@ kernel void main0(constant uniformVelocity& _48 [[buffer(0)]], constant uniformD int j1i_2 = int(j1_2); int k0i_2 = int(k0_2); int k1i_2 = int(k1_2); - float4 _797 = float4((s0_2 * ((t0_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i0i_2, j0i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i0i_2, j0i_2, k1i_2)), 0).x))) + (t1_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i0i_2, j1i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i0i_2, j1i_2, k1i_2)), 0).x))))) + (s1_2 * ((t0_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i1i_2, j0i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i1i_2, j0i_2, k1i_2)), 0).x))) + (t1_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i1i_2, j1i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i1i_2, j1i_2, k1i_2)), 0).x)))))); - outVelocity.write(_797, uint3(id)); + float4 _787 = float4((s0_2 * ((t0_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i0i_2, j0i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i0i_2, j0i_2, k1i_2)), 0).x))) + (t1_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i0i_2, j1i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i0i_2, j1i_2, k1i_2)), 0).x))))) + (s1_2 * ((t0_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i1i_2, j0i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i1i_2, j0i_2, k1i_2)), 0).x))) + (t1_2 * ((u0_2 * inVelocityZ.read(uint3(int3(i1i_2, j1i_2, k0i_2)), 0).x) + (u1_2 * inVelocityZ.read(uint3(int3(i1i_2, j1i_2, k1i_2)), 0).x)))))); + outVelocity.write(_787, uint3(id)); } while (false); } } diff --git a/shaders/bin/advect1.comp.spv b/shaders/bin/advect1.comp.spv index 2843f5e..78dfb91 100644 Binary files a/shaders/bin/advect1.comp.spv and b/shaders/bin/advect1.comp.spv differ diff --git a/shaders/bin/advect2.comp.msl b/shaders/bin/advect2.comp.msl index 26e0e11..beba004 100644 --- a/shaders/bin/advect2.comp.msl +++ b/shaders/bin/advect2.comp.msl @@ -31,9 +31,9 @@ kernel void main0(constant uniformDeltaTime& _62 [[buffer(0)]], texture3d do { float N = float(size.x - 2); - float dtx = _62.deltaTime * (N - 2.0); - float dty = _62.deltaTime * (N - 2.0); - float dtz = _62.deltaTime * (N - 2.0); + float dtx = _62.deltaTime * N; + float dty = _62.deltaTime * N; + float dtz = _62.deltaTime * N; float tmp1 = dtx * inVelocityX.read(uint3(id), 0).x; float tmp2 = dty * inVelocityY.read(uint3(id), 0).x; float tmp3 = dtz * inVelocityZ.read(uint3(id), 0).x; @@ -58,8 +58,8 @@ kernel void main0(constant uniformDeltaTime& _62 [[buffer(0)]], texture3d int j1i = int(j1); int k0i = int(k0); int k1i = int(k1); - float4 _302 = float4((s0 * ((t0 * ((u0 * inDensity.read(uint3(int3(i0i, j0i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i0i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inDensity.read(uint3(int3(i0i, j1i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i0i, j1i, k1i)), 0).x))))) + (s1 * ((t0 * ((u0 * inDensity.read(uint3(int3(i1i, j0i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i1i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inDensity.read(uint3(int3(i1i, j1i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i1i, j1i, k1i)), 0).x)))))); - outDensity.write(_302, uint3(id)); + float4 _298 = float4((s0 * ((t0 * ((u0 * inDensity.read(uint3(int3(i0i, j0i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i0i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inDensity.read(uint3(int3(i0i, j1i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i0i, j1i, k1i)), 0).x))))) + (s1 * ((t0 * ((u0 * inDensity.read(uint3(int3(i1i, j0i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i1i, j0i, k1i)), 0).x))) + (t1 * ((u0 * inDensity.read(uint3(int3(i1i, j1i, k0i)), 0).x) + (u1 * inDensity.read(uint3(int3(i1i, j1i, k1i)), 0).x)))))); + outDensity.write(_298, uint3(id)); } while (false); } diff --git a/shaders/bin/advect2.comp.spv b/shaders/bin/advect2.comp.spv index 6dfa42b..8ac5a1d 100644 Binary files a/shaders/bin/advect2.comp.spv and b/shaders/bin/advect2.comp.spv differ diff --git a/shaders/bin/bnd1.comp.spv b/shaders/bin/bnd1.comp.spv index 9004ae9..6f8f205 100644 Binary files a/shaders/bin/bnd1.comp.spv and b/shaders/bin/bnd1.comp.spv differ diff --git a/shaders/bin/bnd2.comp.spv b/shaders/bin/bnd2.comp.spv index 696ac28..dc512a8 100644 Binary files a/shaders/bin/bnd2.comp.spv and b/shaders/bin/bnd2.comp.spv differ diff --git a/shaders/bin/bnd3.comp.spv b/shaders/bin/bnd3.comp.spv index df356a7..c8bec54 100644 Binary files a/shaders/bin/bnd3.comp.spv and b/shaders/bin/bnd3.comp.spv differ diff --git a/shaders/bin/bnd4.comp.spv b/shaders/bin/bnd4.comp.spv index e023eaf..5ee9085 100644 Binary files a/shaders/bin/bnd4.comp.spv and b/shaders/bin/bnd4.comp.spv differ diff --git a/shaders/bin/bnd5.comp.spv b/shaders/bin/bnd5.comp.spv index 9e1e25e..dcd9889 100644 Binary files a/shaders/bin/bnd5.comp.spv and b/shaders/bin/bnd5.comp.spv differ diff --git a/shaders/bin/composite.comp.msl b/shaders/bin/composite.comp.msl index 1aba1ff..3cb973a 100644 --- a/shaders/bin/composite.comp.msl +++ b/shaders/bin/composite.comp.msl @@ -32,46 +32,36 @@ float3 GetRayDirection(thread const float4x4& inverseView, thread const float4x4 return fast::normalize(worldRay.xyz); } -kernel void main0(constant uniformInverseView& _120 [[buffer(0)]], constant uniformInverseProj& _123 [[buffer(1)]], constant uniformCameraPosition& _149 [[buffer(2)]], texture3d inVelocityX [[texture(0)]], texture3d inVelocityY [[texture(1)]], texture3d inVelocityZ [[texture(2)]], texture3d inDensity [[texture(3)]], texture2d outColor [[texture(4)]], sampler inVelocityXSmplr [[sampler(0)]], sampler inVelocityYSmplr [[sampler(1)]], sampler inVelocityZSmplr [[sampler(2)]], sampler inDensitySmplr [[sampler(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant uniformInverseView& _109 [[buffer(0)]], constant uniformInverseProj& _112 [[buffer(1)]], constant uniformCameraPosition& _138 [[buffer(2)]], texture3d inVelocityX [[texture(0)]], texture3d inVelocityY [[texture(1)]], texture3d inVelocityZ [[texture(2)]], texture3d inDensity [[texture(3)]], texture2d outColor [[texture(4)]], sampler inVelocityXSmplr [[sampler(0)]], sampler inVelocityYSmplr [[sampler(1)]], sampler inVelocityZSmplr [[sampler(2)]], sampler inDensitySmplr [[sampler(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { int3 size = int3(inVelocityX.get_width(), inVelocityX.get_height(), inVelocityX.get_depth()); int2 _imageSize = int2(outColor.get_width(), outColor.get_height()); int2 id = int2(gl_GlobalInvocationID.xy); - bool _93 = any(id >= (_imageSize - int2(1))); - bool _101; - if (!_93) - { - _101 = any(id <= int2(0)); - } - else - { - _101 = _93; - } - if (_101) + if (any(id >= _imageSize)) { return; } float2 texcoord = float2(id) / float2(_imageSize); texcoord.y = 1.0 - texcoord.y; - float4x4 param = _120.inverseView; - float4x4 param_1 = _123.inverseProj; + float4x4 param = _109.inverseView; + float4x4 param_1 = _112.inverseProj; float2 param_2 = texcoord; float3 rayDirection = GetRayDirection(param, param_1, param_2); float4 result = float4(0.0); for (int i = 0; i < 512; i++) { - int3 cell = int3(_149.cameraPosition + ((rayDirection * float(i)) * 1.0)); - bool _164 = any(cell >= size); - bool _172; - if (!_164) + int3 cell = int3(_138.cameraPosition + ((rayDirection * float(i)) * 1.0)); + bool _153 = any(cell >= size); + bool _161; + if (!_153) { - _172 = any(cell < int3(0)); + _161 = any(cell < int3(0)); } else { - _172 = _164; + _161 = _153; } - if (_172) + if (_161) { continue; } diff --git a/shaders/bin/composite.comp.spv b/shaders/bin/composite.comp.spv index 18982d5..a5a1de2 100644 Binary files a/shaders/bin/composite.comp.spv and b/shaders/bin/composite.comp.spv differ diff --git a/shaders/bin/diffuse.comp.json b/shaders/bin/diffuse.comp.json index da4e0d1..645dc26 100644 --- a/shaders/bin/diffuse.comp.json +++ b/shaders/bin/diffuse.comp.json @@ -1 +1 @@ -{ "samplers": 1, "readonly_storage_textures": 0, "readonly_storage_buffers": 0, "readwrite_storage_textures": 1, "readwrite_storage_buffers": 0, "uniform_buffers": 2, "threadcount_x": 8, "threadcount_y": 8, "threadcount_z": 8 } +{ "samplers": 2, "readonly_storage_textures": 0, "readonly_storage_buffers": 0, "readwrite_storage_textures": 1, "readwrite_storage_buffers": 0, "uniform_buffers": 2, "threadcount_x": 8, "threadcount_y": 8, "threadcount_z": 8 } diff --git a/shaders/bin/diffuse.comp.msl b/shaders/bin/diffuse.comp.msl index 1f1bb9e..dcf704b 100644 --- a/shaders/bin/diffuse.comp.msl +++ b/shaders/bin/diffuse.comp.msl @@ -15,7 +15,7 @@ struct uniformDiffusion constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u); -kernel void main0(constant uniformDeltaTime& _55 [[buffer(0)]], constant uniformDiffusion& _61 [[buffer(1)]], texture3d inImage [[texture(0)]], texture3d outImage [[texture(1)]], sampler inImageSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant uniformDeltaTime& _55 [[buffer(0)]], constant uniformDiffusion& _61 [[buffer(1)]], texture3d inSource [[texture(0)]], texture3d inImage [[texture(1)]], texture3d outImage [[texture(2)]], sampler inSourceSmplr [[sampler(0)]], sampler inImageSmplr [[sampler(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { int3 size = int3(inImage.get_width(), inImage.get_height(), inImage.get_depth()); int3 id = int3(gl_GlobalInvocationID); @@ -38,7 +38,7 @@ kernel void main0(constant uniformDeltaTime& _55 [[buffer(0)]], constant uniform float c = 1.0 + (6.0 * a); do { - outImage.write(float4((inImage.read(uint3(id), 0).x + (a * (((((inImage.read(uint3((id + int3(1, 0, 0))), 0).x + inImage.read(uint3((id + int3(-1, 0, 0))), 0).x) + inImage.read(uint3((id + int3(0, 1, 0))), 0).x) + inImage.read(uint3((id + int3(0, -1, 0))), 0).x) + inImage.read(uint3((id + int3(0, 0, 1))), 0).x) + inImage.read(uint3((id + int3(0, 0, -1))), 0).x))) / c), uint3(id)); + outImage.write(float4((inSource.read(uint3(id), 0).x + (a * (((((inImage.read(uint3((id + int3(1, 0, 0))), 0).x + inImage.read(uint3((id + int3(-1, 0, 0))), 0).x) + inImage.read(uint3((id + int3(0, 1, 0))), 0).x) + inImage.read(uint3((id + int3(0, -1, 0))), 0).x) + inImage.read(uint3((id + int3(0, 0, 1))), 0).x) + inImage.read(uint3((id + int3(0, 0, -1))), 0).x))) / c), uint3(id)); } while (false); } diff --git a/shaders/bin/diffuse.comp.spv b/shaders/bin/diffuse.comp.spv index 6d4674f..f9a554a 100644 Binary files a/shaders/bin/diffuse.comp.spv and b/shaders/bin/diffuse.comp.spv differ diff --git a/shaders/bin/project1.comp.spv b/shaders/bin/project1.comp.spv index 04c506e..ba7a94e 100644 Binary files a/shaders/bin/project1.comp.spv and b/shaders/bin/project1.comp.spv differ diff --git a/shaders/bin/project2.comp.json b/shaders/bin/project2.comp.json index 73163ba..0d1932d 100644 --- a/shaders/bin/project2.comp.json +++ b/shaders/bin/project2.comp.json @@ -1 +1 @@ -{ "samplers": 1, "readonly_storage_textures": 0, "readonly_storage_buffers": 0, "readwrite_storage_textures": 1, "readwrite_storage_buffers": 0, "uniform_buffers": 0, "threadcount_x": 8, "threadcount_y": 8, "threadcount_z": 8 } +{ "samplers": 2, "readonly_storage_textures": 0, "readonly_storage_buffers": 0, "readwrite_storage_textures": 1, "readwrite_storage_buffers": 0, "uniform_buffers": 0, "threadcount_x": 8, "threadcount_y": 8, "threadcount_z": 8 } diff --git a/shaders/bin/project2.comp.msl b/shaders/bin/project2.comp.msl index cb62943..cff56b4 100644 --- a/shaders/bin/project2.comp.msl +++ b/shaders/bin/project2.comp.msl @@ -5,7 +5,7 @@ using namespace metal; constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(8u); -kernel void main0(texture3d inDivergence [[texture(0)]], texture3d outPressure [[texture(1)]], sampler inDivergenceSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(texture3d inDivergence [[texture(0)]], texture3d inPressure [[texture(1)]], texture3d outPressure [[texture(2)]], sampler inDivergenceSmplr [[sampler(0)]], sampler inPressureSmplr [[sampler(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { int3 size = int3(inDivergence.get_width(), inDivergence.get_height(), inDivergence.get_depth()); int3 id = int3(gl_GlobalInvocationID); @@ -25,7 +25,7 @@ kernel void main0(texture3d inDivergence [[texture(0)]], texture3d inImage [[texture(0)]], texture2d outColor [[texture(1)]], sampler inImageSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +kernel void main0(constant uniformInverseView& _109 [[buffer(0)]], constant uniformInverseProj& _112 [[buffer(1)]], constant uniformCameraPosition& _138 [[buffer(2)]], texture3d inImage [[texture(0)]], texture2d outColor [[texture(1)]], sampler inImageSmplr [[sampler(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) { int3 size = int3(inImage.get_width(), inImage.get_height(), inImage.get_depth()); int2 _imageSize = int2(outColor.get_width(), outColor.get_height()); int2 id = int2(gl_GlobalInvocationID.xy); - bool _93 = any(id >= (_imageSize - int2(1))); - bool _101; - if (!_93) - { - _101 = any(id <= int2(0)); - } - else - { - _101 = _93; - } - if (_101) + if (any(id >= _imageSize)) { return; } float2 texcoord = float2(id) / float2(_imageSize); texcoord.y = 1.0 - texcoord.y; - float4x4 param = _120.inverseView; - float4x4 param_1 = _123.inverseProj; + float4x4 param = _109.inverseView; + float4x4 param_1 = _112.inverseProj; float2 param_2 = texcoord; float3 rayDirection = GetRayDirection(param, param_1, param_2); float4 result = float4(0.0); for (int i = 0; i < 512; i++) { - int3 cell = int3(_149.cameraPosition + ((rayDirection * float(i)) * 1.0)); - bool _164 = any(cell >= size); - bool _172; - if (!_164) + int3 cell = int3(_138.cameraPosition + ((rayDirection * float(i)) * 1.0)); + bool _153 = any(cell >= size); + bool _161; + if (!_153) { - _172 = any(cell < int3(0)); + _161 = any(cell < int3(0)); } else { - _172 = _164; + _161 = _153; } - if (_172) + if (_161) { continue; } diff --git a/shaders/bin/single.comp.spv b/shaders/bin/single.comp.spv index 6b851c1..460ae72 100644 Binary files a/shaders/bin/single.comp.spv and b/shaders/bin/single.comp.spv differ diff --git a/shaders/bnd1.comp b/shaders/bnd1.comp index baf2e43..6da8b77 100644 --- a/shaders/bnd1.comp +++ b/shaders/bnd1.comp @@ -4,7 +4,7 @@ layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = 1) in; layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outImage; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outImage; layout(set = 2, binding = 0) uniform uniformType { uint type; diff --git a/shaders/bnd2.comp b/shaders/bnd2.comp index c16d19a..ad9d1c6 100644 --- a/shaders/bnd2.comp +++ b/shaders/bnd2.comp @@ -4,7 +4,7 @@ layout(local_size_x = THREADS, local_size_y = 1, local_size_z = THREADS) in; layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outImage; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outImage; layout(set = 2, binding = 0) uniform uniformType { uint type; diff --git a/shaders/bnd3.comp b/shaders/bnd3.comp index d6ad88f..7b4d2ab 100644 --- a/shaders/bnd3.comp +++ b/shaders/bnd3.comp @@ -4,7 +4,7 @@ layout(local_size_x = 1, local_size_y = THREADS, local_size_z = THREADS) in; layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outImage; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outImage; layout(set = 2, binding = 0) uniform uniformType { uint type; diff --git a/shaders/bnd4.comp b/shaders/bnd4.comp index 43370bf..b1605ba 100644 --- a/shaders/bnd4.comp +++ b/shaders/bnd4.comp @@ -4,7 +4,7 @@ layout(local_size_x = 8, local_size_y = 1, local_size_z = 1) in; layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outImage; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outImage; void main() { diff --git a/shaders/bnd5.comp b/shaders/bnd5.comp index d051612..8b5e023 100644 --- a/shaders/bnd5.comp +++ b/shaders/bnd5.comp @@ -4,7 +4,7 @@ layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) in; layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outImage; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outImage; void main() { diff --git a/shaders/composite.comp b/shaders/composite.comp index b7011a1..be302b6 100644 --- a/shaders/composite.comp +++ b/shaders/composite.comp @@ -8,7 +8,7 @@ layout(set = 0, binding = 0) uniform sampler3D inVelocityX; layout(set = 0, binding = 1) uniform sampler3D inVelocityY; layout(set = 0, binding = 2) uniform sampler3D inVelocityZ; layout(set = 0, binding = 3) uniform sampler3D inDensity; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image2D outColor; +layout(set = 1, binding = 0, rgba8) uniform writeonly image2D outColor; layout(set = 2, binding = 0) uniform uniformInverseView { mat4 inverseView; @@ -27,7 +27,7 @@ void main() ivec3 size = textureSize(inVelocityX, 0); ivec2 imageSize = imageSize(outColor); ivec2 id = ivec2(gl_GlobalInvocationID.xy); - if (any(greaterThanEqual(id, imageSize - 1)) || any(lessThanEqual(id, ivec2(0)))) + if (any(greaterThanEqual(id, imageSize))) { return; } diff --git a/shaders/diffuse.comp b/shaders/diffuse.comp index 2d94ace..09c7ac2 100644 --- a/shaders/diffuse.comp +++ b/shaders/diffuse.comp @@ -4,8 +4,9 @@ #include "shader.glsl" layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) in; -layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outImage; +layout(set = 0, binding = 0) uniform sampler3D inSource; +layout(set = 0, binding = 1) uniform sampler3D inImage; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outImage; layout(set = 2, binding = 0) uniform uniformDeltaTime { float deltaTime; @@ -26,5 +27,5 @@ void main() int N = size.x; float a = deltaTime * diffusion * (N - 2) * (N - 2); float c = 1 + 6 * a; - LIN_SOLVE(id, inImage, outImage, a, c); + LIN_SOLVE(id, inSource, inImage, outImage, a, c); } \ No newline at end of file diff --git a/shaders/project1.comp b/shaders/project1.comp index 2c94289..2265a0b 100644 --- a/shaders/project1.comp +++ b/shaders/project1.comp @@ -6,8 +6,8 @@ layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) i layout(set = 0, binding = 0) uniform sampler3D inVelocityX; layout(set = 0, binding = 1) uniform sampler3D inVelocityY; layout(set = 0, binding = 2) uniform sampler3D inVelocityZ; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outPressure; -layout(set = 1, binding = 1, rgba32f) uniform writeonly image3D outDivergence; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outPressure; +layout(set = 1, binding = 1, r32f) uniform writeonly image3D outDivergence; void main() { diff --git a/shaders/project2.comp b/shaders/project2.comp index 70e8e84..66bc7e1 100644 --- a/shaders/project2.comp +++ b/shaders/project2.comp @@ -5,7 +5,8 @@ layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) in; layout(set = 0, binding = 0) uniform sampler3D inDivergence; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outPressure; +layout(set = 0, binding = 1) uniform sampler3D inPressure; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outPressure; void main() { @@ -15,5 +16,5 @@ void main() { return; } - LIN_SOLVE(id, inDivergence, outPressure, 1, 6); + LIN_SOLVE(id, inDivergence, inPressure, outPressure, 1, 6); } \ No newline at end of file diff --git a/shaders/project3.comp b/shaders/project3.comp index ddbda6d..c6ac906 100644 --- a/shaders/project3.comp +++ b/shaders/project3.comp @@ -7,9 +7,9 @@ layout(set = 0, binding = 0) uniform sampler3D inPressure; layout(set = 0, binding = 1) uniform sampler3D inVelocityX; layout(set = 0, binding = 2) uniform sampler3D inVelocityY; layout(set = 0, binding = 3) uniform sampler3D inVelocityZ; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image3D outVelocityX; -layout(set = 1, binding = 1, rgba32f) uniform writeonly image3D outVelocityY; -layout(set = 1, binding = 2, rgba32f) uniform writeonly image3D outVelocityZ; +layout(set = 1, binding = 0, r32f) uniform writeonly image3D outVelocityX; +layout(set = 1, binding = 1, r32f) uniform writeonly image3D outVelocityY; +layout(set = 1, binding = 2, r32f) uniform writeonly image3D outVelocityZ; void main() { diff --git a/shaders/shader.glsl b/shaders/shader.glsl index 2c1aba3..5da3045 100644 --- a/shaders/shader.glsl +++ b/shaders/shader.glsl @@ -14,17 +14,17 @@ vec3 GetRayDirection(mat4 inverseView, mat4 inverseProj, vec2 texcoord) return normalize(worldRay.xyz); } -#define LIN_SOLVE(id, inImage, outImage, a, c) \ +#define LIN_SOLVE(id, srcImage, iterImage, outImage, a, c) \ do \ { \ imageStore(outImage, id, vec4( \ - (texelFetch(inImage, id, 0).x + \ - a * (texelFetch(inImage, id + ivec3( 1, 0, 0 ), 0).x + \ - texelFetch(inImage, id + ivec3(-1, 0, 0 ), 0).x + \ - texelFetch(inImage, id + ivec3( 0, 1, 0 ), 0).x + \ - texelFetch(inImage, id + ivec3( 0,-1, 0 ), 0).x + \ - texelFetch(inImage, id + ivec3( 0, 0, 1 ), 0).x + \ - texelFetch(inImage, id + ivec3( 0, 0,-1 ), 0).x)) / c)); \ + (texelFetch(srcImage, id, 0).x + \ + a * (texelFetch(iterImage, id + ivec3( 1, 0, 0 ), 0).x + \ + texelFetch(iterImage, id + ivec3(-1, 0, 0 ), 0).x + \ + texelFetch(iterImage, id + ivec3( 0, 1, 0 ), 0).x + \ + texelFetch(iterImage, id + ivec3( 0,-1, 0 ), 0).x + \ + texelFetch(iterImage, id + ivec3( 0, 0, 1 ), 0).x + \ + texelFetch(iterImage, id + ivec3( 0, 0,-1 ), 0).x)) / c)); \ } \ while (false) \ @@ -32,9 +32,9 @@ vec3 GetRayDirection(mat4 inverseView, mat4 inverseProj, vec2 texcoord) do \ { \ float N = size.x - 2; \ - float dtx = deltaTime * (N - 2); \ - float dty = deltaTime * (N - 2); \ - float dtz = deltaTime * (N - 2); \ + float dtx = deltaTime * N; \ + float dty = deltaTime * N; \ + float dtz = deltaTime * N; \ float tmp1 = dtx * texelFetch(inVelocityX, id, 0).x; \ float tmp2 = dty * texelFetch(inVelocityY, id, 0).x; \ float tmp3 = dtz * texelFetch(inVelocityZ, id, 0).x; \ diff --git a/shaders/single.comp b/shaders/single.comp index d9ec706..bb5e860 100644 --- a/shaders/single.comp +++ b/shaders/single.comp @@ -5,7 +5,7 @@ layout(local_size_x = THREADS, local_size_y = THREADS, local_size_z = THREADS) in; layout(set = 0, binding = 0) uniform sampler3D inImage; -layout(set = 1, binding = 0, rgba32f) uniform writeonly image2D outColor; +layout(set = 1, binding = 0, rgba8) uniform writeonly image2D outColor; layout(set = 2, binding = 0) uniform uniformInverseView { mat4 inverseView; @@ -24,7 +24,7 @@ void main() ivec3 size = textureSize(inImage, 0); ivec2 imageSize = imageSize(outColor); ivec2 id = ivec2(gl_GlobalInvocationID.xy); - if (any(greaterThanEqual(id, imageSize - 1)) || any(lessThanEqual(id, ivec2(0)))) + if (any(greaterThanEqual(id, imageSize))) { return; } diff --git a/src/main.cpp b/src/main.cpp index a8a89e9..de9af49 100644 --- a/src/main.cpp +++ b/src/main.cpp @@ -67,8 +67,8 @@ struct State { int Size = 128; int Iterations = 5; - float Diffusion = 0.01f; - float Viscosity = 0.01f; + float Diffusion = 0.000004f; + float Viscosity = 0.000004f; std::vector Spawners; NLOHMANN_DEFINE_TYPE_INTRUSIVE(State, Size, Iterations, Diffusion, Viscosity, Spawners) @@ -77,7 +77,6 @@ struct State enum PipelineType { PipelineTypeAdd1, - PipelineTypeAdd2, PipelineTypeClear, PipelineTypeDiffuse, PipelineTypeProject1, @@ -111,6 +110,7 @@ static SDL_GPUTexture* colorTexture; static uint32_t width; static uint32_t height; static ReadWriteTexture textures[TextureTypeCount]; +static SDL_GPUTexture* scratchTexture; static SDL_GPUSampler* sampler; static float speed = 16.0f; static int cooldown; @@ -211,7 +211,6 @@ static bool CreateTextures() bool CreatePipelines() { pipelines[PipelineTypeAdd1] = LoadComputePipeline(device, "add1.comp"); - pipelines[PipelineTypeAdd2] = LoadComputePipeline(device, "add2.comp"); pipelines[PipelineTypeClear] = LoadComputePipeline(device, "clear.comp"); pipelines[PipelineTypeDiffuse] = LoadComputePipeline(device, "diffuse.comp"); pipelines[PipelineTypeProject1] = LoadComputePipeline(device, "project1.comp"); @@ -269,22 +268,6 @@ static void Add1(SDL_GPUCommandBuffer* commandBuffer, TextureType texture, const SDL_EndGPUComputePass(computePass); } -static void Add2(SDL_GPUCommandBuffer* commandBuffer, TextureType texture, float value) -{ - DebugGroup(commandBuffer); - SDL_GPUComputePass* computePass = textures[texture].BeginReadPass(commandBuffer); - if (!computePass) - { - SDL_Log("Failed to begin compute pass: %s", SDL_GetError()); - return; - } - int groups = (state.Size + THREADS - 1) / THREADS; - SDL_BindGPUComputePipeline(computePass, pipelines[PipelineTypeAdd2]); - SDL_PushGPUComputeUniformData(commandBuffer, 0, &value, sizeof(value)); - SDL_DispatchGPUCompute(computePass, groups, groups, groups); - SDL_EndGPUComputePass(computePass); -} - static void Clear(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& texture, float value = 0.0f) { DebugGroup(commandBuffer); @@ -303,6 +286,21 @@ static void Clear(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& texture static bool CreateCells() { + SDL_ReleaseGPUTexture(device, scratchTexture); + SDL_GPUTextureCreateInfo info{}; + info.format = SDL_GPU_TEXTUREFORMAT_R32_FLOAT; + info.type = SDL_GPU_TEXTURETYPE_3D; + info.usage = SDL_GPU_TEXTUREUSAGE_SAMPLER; + info.width = state.Size; + info.height = state.Size; + info.layer_count_or_depth = state.Size; + info.num_levels = 1; + scratchTexture = SDL_CreateGPUTexture(device, &info); + if (!scratchTexture) + { + SDL_Log("Failed to create texture: %s", SDL_GetError()); + return false; + } SDL_GPUCommandBuffer* commandBuffer = SDL_AcquireGPUCommandBuffer(device); if (!commandBuffer) { @@ -365,13 +363,13 @@ static void LoadCallback(void *userdata, const char* const* filelist, int filter try { file >> json; + state = json; } catch (const std::exception& exception) { SDL_Log("Failed to load json: %s, %s", filelist[0], exception.what()); return; } - state = json; CreateCells(); } @@ -391,7 +389,7 @@ static void UpdateSpawners(SDL_GPUCommandBuffer* commandBuffer) { for (int j = 0; j < SDL_arraysize(Spawners); j++) { - bool isSelected = spawner.Texture == j; + bool isSelected = spawner.Texture == Spawners[j]; if (ImGui::Selectable(Textures[Spawners[j]], isSelected)) { spawner.Texture = Spawners[j]; @@ -439,7 +437,7 @@ static void UpdateImGui(SDL_GPUCommandBuffer* commandBuffer) ImGui_ImplSDLGPU3_NewFrame(); ImGui::NewFrame(); ImGui::Begin("Fluid Simulation"); - const char* location = SDL_GetCurrentDirectory(); + char* location = SDL_GetCurrentDirectory(); if (ImGui::Button("Save")) { SDL_ShowSaveFileDialog(SaveCallback, nullptr, window, nullptr, 1, location); @@ -449,6 +447,7 @@ static void UpdateImGui(SDL_GPUCommandBuffer* commandBuffer) { SDL_ShowOpenFileDialog(LoadCallback, nullptr, window, nullptr, 1, location, false); } + SDL_free(location); ImGui::SameLine(); if (ImGui::Button("Reset")) { @@ -457,8 +456,8 @@ static void UpdateImGui(SDL_GPUCommandBuffer* commandBuffer) ImGui::SeparatorText("Settings"); ImGui::SliderFloat("Speed", &speed, 0.0f, 64.0f); ImGui::SliderInt("Iterations", &state.Iterations, 1, 50); - ImGui::SliderFloat("Diffusion", &state.Diffusion, 0.0f, 1.0f); - ImGui::SliderFloat("Viscosity", &state.Viscosity, 0.0f, 1.0f); + ImGui::SliderFloat("Diffusion", &state.Diffusion, 0.0f, 0.0001f, "%.7f", ImGuiSliderFlags_Logarithmic); + ImGui::SliderFloat("Viscosity", &state.Viscosity, 0.0f, 0.0001f, "%.7f", ImGuiSliderFlags_Logarithmic); if (ImGui::SliderInt("Size", &state.Size, 16, 256)) { CreateCells(); @@ -477,7 +476,7 @@ static void UpdateImGui(SDL_GPUCommandBuffer* commandBuffer) ImGui_ImplSDLGPU3_PrepareDrawData(ImGui::GetDrawData(), commandBuffer); } -static void Diffuse(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& texture, float diffusion) +static void Diffuse1(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& texture, float diffusion) { DebugGroup(commandBuffer); SDL_GPUComputePass* computePass = texture.BeginWritePass(commandBuffer); @@ -486,12 +485,14 @@ static void Diffuse(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& textu SDL_Log("Failed to begin compute pass: %s", SDL_GetError()); return; } - SDL_GPUTextureSamplerBinding textureBindings{}; - textureBindings.sampler = sampler; - textureBindings.texture = texture.GetReadTexture(); + SDL_GPUTextureSamplerBinding textureBindings[2]{}; + textureBindings[0].sampler = sampler; + textureBindings[0].texture = scratchTexture; + textureBindings[1].sampler = sampler; + textureBindings[1].texture = texture.GetReadTexture(); int groups = (state.Size + THREADS - 1) / THREADS; SDL_BindGPUComputePipeline(computePass, pipelines[PipelineTypeDiffuse]); - SDL_BindGPUComputeSamplers(computePass, 0, &textureBindings, 1); + SDL_BindGPUComputeSamplers(computePass, 0, textureBindings, 2); SDL_PushGPUComputeUniformData(commandBuffer, 0, &speed, sizeof(speed)); SDL_PushGPUComputeUniformData(commandBuffer, 1, &diffusion, sizeof(diffusion)); SDL_DispatchGPUCompute(computePass, groups, groups, groups); @@ -536,12 +537,14 @@ static void Project2(SDL_GPUCommandBuffer* commandBuffer) SDL_Log("Failed to begin compute pass: %s", SDL_GetError()); return; } - SDL_GPUTextureSamplerBinding textureBinding{}; - textureBinding.sampler = sampler; - textureBinding.texture = textures[TextureTypeDivergence].GetReadTexture(); + SDL_GPUTextureSamplerBinding textureBindings[2]{}; + textureBindings[0].sampler = sampler; + textureBindings[0].texture = textures[TextureTypeDivergence].GetReadTexture(); + textureBindings[1].sampler = sampler; + textureBindings[1].texture = textures[TextureTypePressure].GetReadTexture(); int groups = (state.Size + THREADS - 1) / THREADS; SDL_BindGPUComputePipeline(computePass, pipelines[PipelineTypeProject2]); - SDL_BindGPUComputeSamplers(computePass, 0, &textureBinding, 1); + SDL_BindGPUComputeSamplers(computePass, 0, textureBindings, 2); SDL_DispatchGPUCompute(computePass, groups, groups, groups); SDL_EndGPUComputePass(computePass); textures[TextureTypePressure].Swap(); @@ -739,6 +742,44 @@ static void Bnd(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& texture, texture.Swap(); } +static void Project(SDL_GPUCommandBuffer* commandBuffer) +{ + Project1(commandBuffer); + Bnd(commandBuffer, textures[TextureTypeDivergence], 0); + Bnd(commandBuffer, textures[TextureTypePressure], 0); + for (int i = 0; i < state.Iterations; i++) + { + Project2(commandBuffer); + Bnd(commandBuffer, textures[TextureTypePressure], 0); + } + Project3(commandBuffer); + Bnd(commandBuffer, textures[TextureTypeVelocityX], 1); + Bnd(commandBuffer, textures[TextureTypeVelocityY], 2); + Bnd(commandBuffer, textures[TextureTypeVelocityZ], 3); +} + +static void Diffuse(SDL_GPUCommandBuffer* commandBuffer, ReadWriteTexture& texture, float diffusion, int type) +{ + DebugGroup(commandBuffer); + SDL_GPUCopyPass* copyPass = SDL_BeginGPUCopyPass(commandBuffer); + if (!copyPass) + { + SDL_Log("Failed to begin copy pass: %s", SDL_GetError()); + return; + } + SDL_GPUTextureLocation source{}; + source.texture = texture.GetReadTexture(); + SDL_GPUTextureLocation destination{}; + destination.texture = scratchTexture; + SDL_CopyGPUTextureToTexture(copyPass, &source, &destination, state.Size, state.Size, state.Size, false); + SDL_EndGPUCopyPass(copyPass); + for (int i = 0; i < state.Iterations; i++) + { + Diffuse1(commandBuffer, texture, diffusion); + Bnd(commandBuffer, texture, type); + } +} + static void RenderComposite(SDL_GPUCommandBuffer* commandBuffer) { DebugGroup(commandBuffer); @@ -878,27 +919,10 @@ static void Update() UpdateViewProj(); if (cooldown <= 0) { - for (int i = 0; i < state.Iterations; i++) - { - Diffuse(commandBuffer, textures[TextureTypeVelocityX], state.Viscosity); - Diffuse(commandBuffer, textures[TextureTypeVelocityY], state.Viscosity); - Diffuse(commandBuffer, textures[TextureTypeVelocityZ], state.Viscosity); - Bnd(commandBuffer, textures[TextureTypeVelocityX], 1); - Bnd(commandBuffer, textures[TextureTypeVelocityY], 2); - Bnd(commandBuffer, textures[TextureTypeVelocityZ], 3); - } - Project1(commandBuffer); - Bnd(commandBuffer, textures[TextureTypeDivergence], 0); - Bnd(commandBuffer, textures[TextureTypePressure], 0); - for (int i = 0; i < state.Iterations; i++) - { - Project2(commandBuffer); - Bnd(commandBuffer, textures[TextureTypePressure], 0); - } - Project3(commandBuffer); - Bnd(commandBuffer, textures[TextureTypeVelocityX], 1); - Bnd(commandBuffer, textures[TextureTypeVelocityY], 2); - Bnd(commandBuffer, textures[TextureTypeVelocityZ], 3); + Diffuse(commandBuffer, textures[TextureTypeVelocityX], state.Viscosity, 1); + Diffuse(commandBuffer, textures[TextureTypeVelocityY], state.Viscosity, 2); + Diffuse(commandBuffer, textures[TextureTypeVelocityZ], state.Viscosity, 3); + Project(commandBuffer); Advect1(commandBuffer, TextureTypeVelocityX); Advect1(commandBuffer, TextureTypeVelocityY); Advect1(commandBuffer, TextureTypeVelocityZ); @@ -908,10 +932,8 @@ static void Update() Bnd(commandBuffer, textures[TextureTypeVelocityX], 1); Bnd(commandBuffer, textures[TextureTypeVelocityY], 2); Bnd(commandBuffer, textures[TextureTypeVelocityZ], 3); - Project1(commandBuffer); - Project2(commandBuffer); - Project3(commandBuffer); - Diffuse(commandBuffer, textures[TextureTypeDensity], state.Diffusion); + Project(commandBuffer); + Diffuse(commandBuffer, textures[TextureTypeDensity], state.Diffusion, 0); Advect2(commandBuffer); Bnd(commandBuffer, textures[TextureTypeDensity], 0); cooldown = kCooldown; @@ -1015,12 +1037,12 @@ int main(int argc, char** argv) { textures[i].Free(device); } + SDL_ReleaseGPUTexture(device, scratchTexture); SDL_ReleaseGPUTexture(device, colorTexture); SDL_ReleaseGPUSampler(device, sampler); for (int i = 0; i < PipelineTypeCount; i++) { SDL_ReleaseGPUComputePipeline(device, pipelines[i]); - pipelines[i] = nullptr; } ImGui_ImplSDLGPU3_Shutdown(); ImGui_ImplSDL3_Shutdown();