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..5964754 --- /dev/null +++ b/examples/Fluid2D/Fluid2D.cpp @@ -0,0 +1,577 @@ +#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); + + const auto shader = graphics->CreateShader(data.data(), static_cast(data.size())); + + assert(shader != nullptr); + + return LLGI::CreateSharedPtr(shader); + } + 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); + } + + const auto shader = graphics->CreateShader(data.data(), static_cast(data.size())); + + assert(shader != nullptr); + + return LLGI::CreateSharedPtr(shader); + } +} + +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()); + assert(calcExternalPipeline_->Compile()); + + { + 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()); + assert(buildGridPipeline_->Compile()); + + { + 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()); + assert(bitonicSortPipeline_->Compile()); + + // clearGridIndices + std::cout << "Creating clearGridIndicesPipeline_" << std::endl; + clearGridIndicesPipeline_ = LLGI::CreateSharedPtr(graphics->CreatePiplineState()); + clearGridIndicesPipeline_->SetShader(LLGI::ShaderStageType::Compute, clearGridIndicesShader_.get()); + assert(clearGridIndicesPipeline_->Compile()); + + // 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()); + assert(buildGridIndicesPipeline_->Compile()); + + { + 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()); + assert(calcScalingFactorPipeline_->Compile()); + + { + 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()); + assert(calcCorrectPositionPipeline_->Compile()); + + { + 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()); + assert(integratePipeline_->Compile()); + + { + 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()); + assert(buildVBIBPipeline_->Compile()); + + { + const auto cb = (BuildVBIBCB*)buildVBIBConstant_->Lock(); + + cb->ParticleRadius = ParticleRadius; + cb->Color[0] = 50u; + cb->Color[1] = 50u; + cb->Color[2] = 255u; + cb->Color[3] = 255u; + + cb->FixedColor[0] = 0u; + cb->FixedColor[1] = 0u; + cb->FixedColor[2] = 0u; + cb->FixedColor[3] = 255u; + + 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(), 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); + 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(); + + 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[2] = LLGI::VertexLayoutFormat::R32G32_FLOAT; + pip->VertexLayouts[3] = LLGI::VertexLayoutFormat::R32G32_FLOAT; + pip->VertexLayoutNames[0] = "POSITION"; + pip->VertexLayoutNames[1] = "COLOR"; + 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; + + 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()); + assert(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..d301fc3 --- /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[4]; + 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 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/buildVBIB.comp b/examples/Fluid2D/Shaders/GLSL_GL/buildVBIB.comp new file mode 100644 index 0000000..79977c1 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_GL/buildVBIB.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_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/bitonicSort.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/bitonicSort.comp new file mode 100644 index 0000000..c5ce564 --- /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 = 2, binding = 0, 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..f2d6fbe --- /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 = 2, binding = 1, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +layout(set = 2, 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(_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..6416368 --- /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 = 2, binding = 0, std430) buffer gridTable +{ + ivec2 _data[]; +} gridTable_1; + +layout(set = 2, binding = 1, 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/buildVBIB.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.comp new file mode 100644 index 0000000..fe26b15 --- /dev/null +++ b/examples/Fluid2D/Shaders/GLSL_VULKAN/buildVBIB.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 = 2, binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 2, binding = 1, std430) buffer vertex_ +{ + Vertex _data[]; +} vertex_1; + +layout(set = 2, 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(_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..a023e3c --- /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 = 2, binding = 2, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +layout(set = 2, binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 2, 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 = _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..2acb5fe --- /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 = 2, binding = 0, 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..4fc481b --- /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 = 2, binding = 2, std430) buffer gridIndicesTable +{ + ivec2 _data[]; +} gridIndicesTable_1; + +layout(set = 2, binding = 0, std430) buffer particles +{ + Particle _data[]; +} particles_1; + +layout(set = 2, 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 = _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..96f27da --- /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 = 2, 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_VULKAN/integrate.comp b/examples/Fluid2D/Shaders/GLSL_VULKAN/integrate.comp new file mode 100644 index 0000000..ab74f5e --- /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 = 2, binding = 0, 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/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/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/buildVBIB.comp b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp new file mode 100644 index 0000000..859f13e --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/buildVBIB.comp @@ -0,0 +1,64 @@ +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; + float4 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); + +[numthreads(1, 1, 1)] +void main(uint3 dtid : SV_DispatchThreadID) +{ + 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.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); + 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..5762e4e --- /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 diff --git a/examples/Fluid2D/Shaders/HLSL_DX12/render.vert b/examples/Fluid2D/Shaders/HLSL_DX12/render.vert new file mode 100644 index 0000000..2ea7eeb --- /dev/null +++ b/examples/Fluid2D/Shaders/HLSL_DX12/render.vert @@ -0,0 +1,23 @@ +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 / 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; + return output; +} \ No newline at end of file diff --git a/examples/Fluid2D/Shaders/Metal/bitonicSort.comp b/examples/Fluid2D/Shaders/Metal/bitonicSort.comp new file mode 100644 index 0000000..2c7711d --- /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(10)]], 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..28b79d9 --- /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(10)]], device gridTable& gridTable_1 [[buffer(11)]], 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..8f29df8 --- /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(10)]], device gridTable& gridIndicesTable [[buffer(11)]], 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/buildVBIB.comp b/examples/Fluid2D/Shaders/Metal/buildVBIB.comp new file mode 100644 index 0000000..52a02e7 --- /dev/null +++ b/examples/Fluid2D/Shaders/Metal/buildVBIB.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(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; + _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..3efbc2b --- /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(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; + _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..28b0408 --- /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(10)]], 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..3d7f579 --- /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(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; + _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..4d23d2d --- /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(10)]], 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..626907d --- /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(10)]], 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/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/bitonicSort.comp.spv b/examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv new file mode 100644 index 0000000..886fe47 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/bitonicSort.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/buildGrid.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildGrid.comp.spv new file mode 100644 index 0000000..f1a2357 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/buildGrid.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/buildGridIndices.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildGridIndices.comp.spv new file mode 100644 index 0000000..4ed858f Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/buildGridIndices.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv b/examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv new file mode 100644 index 0000000..26bec3d Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/buildVBIB.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/calcCorrectPosition.comp.spv b/examples/Fluid2D/Shaders/SPIRV/calcCorrectPosition.comp.spv new file mode 100644 index 0000000..b0e838b Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/calcCorrectPosition.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/calcExternalForce.comp.spv b/examples/Fluid2D/Shaders/SPIRV/calcExternalForce.comp.spv new file mode 100644 index 0000000..7e2d983 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/calcExternalForce.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/calcScalingFactor.comp.spv b/examples/Fluid2D/Shaders/SPIRV/calcScalingFactor.comp.spv new file mode 100644 index 0000000..e68ed00 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/calcScalingFactor.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/clearGridIndices.comp.spv b/examples/Fluid2D/Shaders/SPIRV/clearGridIndices.comp.spv new file mode 100644 index 0000000..e31c1ee Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/clearGridIndices.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv b/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv new file mode 100644 index 0000000..bc42d78 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/integrate.comp.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/render.frag.spv b/examples/Fluid2D/Shaders/SPIRV/render.frag.spv new file mode 100644 index 0000000..f7db3c1 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/render.frag.spv differ diff --git a/examples/Fluid2D/Shaders/SPIRV/render.vert.spv b/examples/Fluid2D/Shaders/SPIRV/render.vert.spv new file mode 100644 index 0000000..804af04 Binary files /dev/null and b/examples/Fluid2D/Shaders/SPIRV/render.vert.spv differ 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); }