benjamin
/
VulkanSimulation
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

soft edges

Browse Source
feature/softbody-runtime-control
Benjamin Kraft 4 months ago
parent c462321a41
commit 246408af8d
14 changed files with 7078 additions and 148 deletions
Whitespace
Show all changes Ignore whitespace when comparing lines Ignore changes in amount of whitespace Ignore changes in whitespace at EOL
Unified View
Diff Options
Show Stats Download Patch File Download Diff File
  1. 3
      CMakeLists.txt
  2. 47
      include/constraints.hpp
  3. 25
      include/simulation.hpp
  4. 19
      include/soft_body.hpp
  5. 2
      include/tetgen.h
  6. 3
      include/vulkan/vertex.hpp
  7. 6009
      models/bunny_medium.ply
  8. 598
      models/sphere_medium.ply
  9. 101
      shaders/normal.comp
  10. 109
      shaders/pbd.comp
  11. 22
      src/constraints.cpp
  12. 152
      src/simulation.cpp
  13. 130
      src/soft_body.cpp
  14. 2
      src/vulkan/descriptor_pool.cpp

3
CMakeLists.txt
Unescape View File

@ -7,6 +7,7 @@ file(GLOB_RECURSE SRC_FILES src/**.cpp)
add_executable(VulkanSimulation ${SRC_FILES}) add_executable(VulkanSimulation ${SRC_FILES})
find_package(Vulkan REQUIRED) find_package(Vulkan REQUIRED)
find_package(OpenMP REQUIRED)
include(deps.cmake) include(deps.cmake)
if (WIN32) if (WIN32)
@ -18,7 +19,7 @@ FetchContent_MakeAvailable(tetgen)
FetchContent_MakeAvailable(assimp) FetchContent_MakeAvailable(assimp)
target_compile_options(assimp PRIVATE -Wno-unknown-pragmas) target_compile_options(assimp PRIVATE -Wno-unknown-pragmas)
target_link_libraries(VulkanSimulation glm glfw Vulkan::Vulkan GPUOpen::VulkanMemoryAllocator assimp tet) target_link_libraries(VulkanSimulation glm glfw Vulkan::Vulkan GPUOpen::VulkanMemoryAllocator assimp tet OpenMP::OpenMP_CXX)
target_include_directories(VulkanSimulation PRIVATE include) target_include_directories(VulkanSimulation PRIVATE include)
target_compile_definitions(VulkanSimulation PRIVATE target_compile_definitions(VulkanSimulation PRIVATE
GLM_FORCE_RADIANS GLM_FORCE_RADIANS

47
include/constraints.hpp
Unescape View File

@ -1,6 +1,10 @@
#pragma once #pragma once
#include <cstdint> #include <cstdint>
#include <algorithm>
#include <vector>
using std::vector;
struct Edge { struct Edge {
uint32_t a; uint32_t a;
@ -25,3 +29,46 @@ struct Tetrahedron {
uint32_t d; uint32_t d;
float volume; float volume;
}; };
struct ConstraintData {
vector<Edge> edges;
vector<Triangle> triangles;
vector<Tetrahedron> tetrahedra;
struct Partition {
uint32_t offset;
uint32_t size;
};
uint32_t partitionCount;
vector<Partition> edgePartitions;
vector<Partition> trianglePartitions;
vector<Partition> tetrahedronPartitions;
void recordNewPartition();
void writePartitionInformation();
private:
uint32_t prePartitionEdgeCount;
uint32_t prePartitionTetrahedronCount;
};
struct Constraint {
virtual ~Constraint() {}
virtual void writeData(ConstraintData& dataLists) const {};
};
struct DistanceConstraint : Edge, Constraint {
explicit DistanceConstraint(const Edge &edge) : Edge(edge) {}
void writeData(ConstraintData &dataLists) const override;
};
struct AreaConstraint : Triangle, Constraint {
explicit AreaConstraint(const Triangle& triangle) : Triangle(triangle) {}
void writeData(ConstraintData &dataLists) const override;
};
struct VolumeConstraint : Tetrahedron, Constraint {
explicit VolumeConstraint(const Tetrahedron& tetrahedron) : Tetrahedron(tetrahedron) {}
void writeData(ConstraintData &dataLists) const override;
};

25
include/simulation.hpp
Unescape View File

@ -1,6 +1,8 @@
#pragma once #pragma once
#include "application.hpp" #include "application.hpp"
#include "constraints.hpp"
#include <glm/vec3.hpp>
class SoftBody; class SoftBody;
class Buffer; class Buffer;
@ -17,8 +19,31 @@ private:
unique_ptr<Buffer> tetrahedronBuffer; unique_ptr<Buffer> tetrahedronBuffer;
void createMeshBuffers(); void createMeshBuffers();
struct SizeInformation {
uint32_t vertexCount;
uint32_t faceCount;
uint32_t edgeCount;
uint32_t triangleCount;
uint32_t tetrahedronCount;
};
unique_ptr<Buffer> sizeInformationBuffer; unique_ptr<Buffer> sizeInformationBuffer;
struct Properties {
glm::vec3 gravity;
float dt;
uint32_t k;
};
unique_ptr<Buffer> propertiesBuffer;
Properties properties {};
struct PBDPushData {
uint32_t state;
ConstraintData::Partition edgePartition;
ConstraintData::Partition tetrahedronPartition;
};
ConstraintData constraintData {};
unique_ptr<ComputePipeline> pbdPipeline; unique_ptr<ComputePipeline> pbdPipeline;
unique_ptr<ComputePipeline> normalPipeline; unique_ptr<ComputePipeline> normalPipeline;
void createComputePipelines(); void createComputePipelines();

19
include/soft_body.hpp
Unescape View File

@ -4,14 +4,17 @@
#include <string> #include <string>
#include <vector> #include <vector>
#include <glm/vec3.hpp> #include <glm/vec3.hpp>
#include <unordered_map>
#include <unordered_set>
#include "constraints.hpp"
using std::unordered_map;
using std::unordered_set;
using std::vector; using std::vector;
struct Vertex; struct Vertex;
struct Edge;
struct Triangle; typedef unordered_map<const Constraint*, vector<const Constraint *>> Graph;
struct Face;
struct Tetrahedron;
class Mesh; class Mesh;
@ -21,18 +24,16 @@ public:
uint32_t firstIndex = 0; uint32_t firstIndex = 0;
int32_t vertexOffset = 0; int32_t vertexOffset = 0;
uint32_t partitionCount = 0;
float compliance; float compliance;
vector<Vertex> vertices; vector<Vertex> vertices;
vector<Edge> edges;
vector<Triangle> triangles;
vector<Face> faces; vector<Face> faces;
vector<Tetrahedron> tetrahedra; ConstraintData constraintData;
void applyVertexOffset(const glm::vec3& offset); void applyVertexOffset(const glm::vec3& offset);
void applyIndexOffset(int32_t offset);
SoftBody& operator =(const SoftBody& other) = delete; SoftBody& operator =(const SoftBody& other) = delete;
private: private:
void splitConstraints();
void reorderConstraintIndices(const vector<unordered_set<const Constraint *>> &partitions);
}; };

2
include/tetgen.h
Unescape View File

@ -1000,7 +1000,7 @@ public:
// // // //
// Each arraypool contains an array of pointers to a number of blocks. Each // // Each arraypool contains an array of pointers to a number of blocks. Each //
// block contains the same fixed number of objects. Each index of the array // // block contains the same fixed number of objects. Each index of the array //
// addresses a particular object in the pool. The most significant bits add- // // addresses a particular object in the pool. The most significant bits writeData- //
// ress the index of the block containing the object. The less significant // // ress the index of the block containing the object. The less significant //
// bits address this object within the block. // // bits address this object within the block. //
// // // //

3
include/vulkan/vertex.hpp
Unescape View File

@ -8,6 +8,9 @@ struct Vertex {
alignas(16) glm::vec3 position; alignas(16) glm::vec3 position;
alignas(16) glm::vec3 color; alignas(16) glm::vec3 color;
alignas(16) glm::vec3 normal; alignas(16) glm::vec3 normal;
alignas(16) glm::vec3 velocity;
glm::vec3 prevPosition;
float inverseMass;
static VkVertexInputBindingDescription getBindingDescription(); static VkVertexInputBindingDescription getBindingDescription();

6009
models/bunny_medium.ply
View File

File diff suppressed because it is too large Load Diff

598
models/sphere_medium.ply
Unescape View File

@ -0,0 +1,598 @@
ply
format ascii 1.0
comment Created by Blender 3.0.1 - www.blender.org
element vertex 310
property float x
property float y
property float z
property float s
property float t
element face 276
property list uchar uint vertex_indices
end_header
0.000000 0.707107 0.707107 0.739130 0.750000
0.000000 0.866025 0.500000 0.739130 0.833333
-0.134898 0.866025 0.481459 0.695652 0.833333
-0.190775 0.707107 0.680885 0.695652 0.750000
-0.000000 -0.866025 0.500000 0.739130 0.166667
0.000000 -0.707107 0.707107 0.739130 0.250000
-0.190775 -0.707107 0.680885 0.695652 0.250000
-0.134898 -0.866025 0.481459 0.695652 0.166667
0.000000 -0.258819 0.965926 0.739130 0.416667
0.000000 -0.000000 1.000000 0.739130 0.500000
-0.269797 -0.000000 0.962917 0.695652 0.500000
-0.260604 -0.258819 0.930107 0.695652 0.416667
0.000000 0.500000 0.866025 0.739130 0.666667
-0.233651 0.500000 0.833911 0.695652 0.666667
-0.000000 -0.965926 0.258819 0.739130 0.083333
-0.069829 -0.965926 0.249221 0.695652 0.083333
0.000000 0.965926 0.258819 0.739130 0.916667
0.000000 1.000000 -0.000000 0.717391 1.000000
-0.069829 0.965926 0.249221 0.695652 0.916667
0.000000 -0.500000 0.866025 0.739130 0.333333
-0.233651 -0.500000 0.833911 0.695652 0.333333
0.000000 0.258819 0.965926 0.739130 0.583333
-0.260604 0.258819 0.930107 0.695652 0.583333
0.000000 -1.000000 -0.000000 0.717391 0.000000
-0.501880 -0.258819 0.825306 0.652174 0.416667
-0.449973 -0.500000 0.739949 0.652174 0.333333
-0.449973 0.500000 0.739949 0.652174 0.666667
-0.501880 0.258819 0.825306 0.652174 0.583333
0.000000 -1.000000 -0.000000 0.673913 0.000000
-0.134478 -0.965926 0.221140 0.652174 0.083333
-0.134478 0.965926 0.221140 0.652174 0.916667
-0.259792 0.866025 0.427210 0.652174 0.833333
-0.367401 -0.707107 0.604166 0.652174 0.250000
-0.519584 -0.000000 0.854419 0.652174 0.500000
-0.367401 0.707107 0.604166 0.652174 0.750000
-0.259792 -0.866025 0.427210 0.652174 0.166667
0.000000 1.000000 -0.000000 0.673913 1.000000
-0.730836 -0.000000 0.682553 0.608696 0.500000
-0.705933 -0.258819 0.659296 0.608696 0.416667
-0.516779 0.707107 0.482638 0.608696 0.750000
-0.632923 0.500000 0.591108 0.608696 0.666667
-0.365418 -0.866025 0.341277 0.608696 0.166667
-0.189154 -0.965926 0.176658 0.608696 0.083333
0.000000 1.000000 -0.000000 0.630435 1.000000
-0.189154 0.965926 0.176658 0.608696 0.916667
-0.632923 -0.500000 0.591108 0.608696 0.333333
-0.705933 0.258819 0.659296 0.608696 0.583333
0.000000 -1.000000 -0.000000 0.630435 0.000000
-0.365418 0.866025 0.341276 0.608696 0.833333
-0.516779 -0.707107 0.482638 0.608696 0.250000
-0.768931 -0.500000 0.398428 0.565217 0.333333
-0.627830 -0.707107 0.325315 0.565217 0.250000
-0.857631 0.258819 0.444389 0.565217 0.583333
-0.887885 -0.000000 0.460065 0.565217 0.500000
-0.443943 0.866025 0.230032 0.565217 0.833333
-0.627830 0.707107 0.325315 0.565217 0.750000
-0.443943 -0.866025 0.230032 0.565217 0.166667
-0.857631 -0.258819 0.444389 0.565217 0.416667
-0.768931 0.500000 0.398428 0.565217 0.666667
-0.229802 -0.965926 0.119074 0.565217 0.083333
0.000000 1.000000 -0.000000 0.586957 1.000000
-0.229802 0.965926 0.119074 0.565217 0.916667
0.000000 -1.000000 -0.000000 0.586957 0.000000
0.000000 1.000000 -0.000000 0.543478 1.000000
-0.253406 0.965926 0.052658 0.521739 0.916667
-0.945723 -0.258819 0.196523 0.521739 0.416667
-0.847912 -0.500000 0.176198 0.521739 0.333333
-0.847912 0.500000 0.176198 0.521739 0.666667
-0.945723 0.258819 0.196523 0.521739 0.583333
0.000000 -1.000000 -0.000000 0.543478 0.000000
-0.253406 -0.965926 0.052658 0.521739 0.083333
-0.489542 0.866025 0.101728 0.521739 0.833333
-0.692317 -0.707107 0.143865 0.521739 0.250000
-0.979084 -0.000000 0.203456 0.521739 0.500000
-0.692317 0.707107 0.143865 0.521739 0.750000
-0.489542 -0.866025 0.101728 0.521739 0.166667
-0.705458 -0.707107 -0.048255 0.478261 0.250000
-0.498835 -0.866025 -0.034121 0.478261 0.166667
-0.997669 -0.000000 -0.068243 0.478261 0.500000
-0.963674 -0.258819 -0.065917 0.478261 0.416667
-0.705458 0.707107 -0.048255 0.478261 0.750000
-0.864007 0.500000 -0.059100 0.478261 0.666667
-0.258216 -0.965926 -0.017663 0.478261 0.083333
0.000000 1.000000 -0.000000 0.500000 1.000000
-0.258216 0.965926 -0.017663 0.478261 0.916667
-0.864006 -0.500000 -0.059100 0.478261 0.333333
-0.963674 0.258819 -0.065917 0.478261 0.583333
0.000000 -1.000000 -0.000000 0.500000 0.000000
-0.498834 0.866025 -0.034121 0.478261 0.833333
-0.243875 0.965926 -0.086673 0.434783 0.916667
-0.471130 0.866025 -0.167440 0.434783 0.833333
-0.816022 -0.500000 -0.290014 0.434783 0.333333
-0.666279 -0.707107 -0.236796 0.434783 0.250000
-0.910154 0.258819 -0.323469 0.434783 0.583333
-0.942261 -0.000000 -0.334880 0.434783 0.500000
-0.666279 0.707107 -0.236796 0.434783 0.750000
-0.471131 -0.866025 -0.167440 0.434783 0.166667
-0.910154 -0.258819 -0.323469 0.434783 0.416667
-0.816022 0.500000 -0.290014 0.434783 0.666667
-0.243875 -0.965926 -0.086673 0.434783 0.083333
0.000000 1.000000 -0.000000 0.456522 1.000000
0.000000 -1.000000 -0.000000 0.456522 0.000000
-0.408485 -0.866025 -0.288340 0.391304 0.166667
-0.211448 -0.965926 -0.149256 0.391304 0.083333
0.000000 1.000000 -0.000000 0.413043 1.000000
-0.211447 0.965926 -0.149256 0.391304 0.916667
-0.789132 -0.258819 -0.557030 0.391304 0.416667
-0.707517 -0.500000 -0.499420 0.391304 0.333333
-0.707517 0.500000 -0.499420 0.391304 0.666667
-0.789132 0.258819 -0.557030 0.391304 0.583333
0.000000 -1.000000 -0.000000 0.413043 0.000000
-0.408485 0.866025 -0.288340 0.391304 0.833333
-0.577685 -0.707107 -0.407775 0.391304 0.250000
-0.816970 -0.000000 -0.576680 0.391304 0.500000
-0.577685 0.707107 -0.407775 0.391304 0.750000
-0.315544 0.866025 -0.387856 0.347826 0.833333
-0.446246 0.707107 -0.548511 0.347826 0.750000
-0.446246 -0.707107 -0.548511 0.347826 0.250000
-0.315544 -0.866025 -0.387856 0.347826 0.166667
-0.631088 -0.000000 -0.775711 0.347826 0.500000
-0.609584 -0.258819 -0.749279 0.347826 0.416667
-0.546538 0.500000 -0.671786 0.347826 0.666667
-0.163338 -0.965926 -0.200769 0.347826 0.083333
0.000000 1.000000 -0.000000 0.369565 1.000000
-0.163338 0.965926 -0.200769 0.347826 0.916667
-0.546538 -0.500000 -0.671786 0.347826 0.333333
-0.609584 0.258819 -0.749279 0.347826 0.583333
0.000000 -1.000000 -0.000000 0.369565 0.000000
0.000000 -1.000000 -0.000000 0.326087 0.000000
-0.103114 -0.965926 -0.237392 0.304348 0.083333
-0.103114 0.965926 -0.237392 0.304348 0.916667
-0.199200 0.866025 -0.458606 0.304348 0.833333
-0.345025 -0.500000 -0.794328 0.304348 0.333333
-0.281712 -0.707107 -0.648566 0.304348 0.250000
-0.384826 0.258819 -0.885958 0.304348 0.583333
-0.398401 -0.000000 -0.917211 0.304348 0.500000
-0.281712 0.707107 -0.648566 0.304348 0.750000
-0.199201 -0.866025 -0.458606 0.304348 0.166667
-0.384826 -0.258819 -0.885958 0.304348 0.416667
-0.345025 0.500000 -0.794328 0.304348 0.666667
0.000000 1.000000 -0.000000 0.326087 1.000000
-0.096284 0.707107 -0.700521 0.260870 0.750000
-0.117924 0.500000 -0.857959 0.260870 0.666667
-0.068083 -0.866025 -0.495343 0.260870 0.166667
-0.035242 -0.965926 -0.256409 0.260870 0.083333
0.000000 1.000000 -0.000000 0.282609 1.000000
-0.035242 0.965926 -0.256409 0.260870 0.916667
-0.131527 -0.258819 -0.956929 0.260870 0.416667
-0.117924 -0.500000 -0.857959 0.260870 0.333333
-0.131527 0.258819 -0.956929 0.260870 0.583333
0.000000 -1.000000 -0.000000 0.282609 0.000000
-0.068083 0.866025 -0.495343 0.260870 0.833333
-0.096284 -0.707107 -0.700521 0.260870 0.250000
-0.136166 -0.000000 -0.990686 0.260870 0.500000
0.068083 0.866025 -0.495343 0.217391 0.833333
0.096285 0.707107 -0.700521 0.217391 0.750000
0.096285 -0.707107 -0.700521 0.217391 0.250000
0.068083 -0.866025 -0.495343 0.217391 0.166667
0.136167 -0.000000 -0.990686 0.217391 0.500000
0.131527 -0.258819 -0.956929 0.217391 0.416667
0.117924 0.500000 -0.857959 0.217391 0.666667
0.035243 -0.965926 -0.256409 0.217391 0.083333
0.000000 1.000000 -0.000000 0.239130 1.000000
0.035243 0.965926 -0.256408 0.217391 0.916667
0.117924 -0.500000 -0.857959 0.217391 0.333333
0.131527 0.258819 -0.956929 0.217391 0.583333
0.000000 -1.000000 -0.000000 0.239130 0.000000
0.345026 0.500000 -0.794328 0.173913 0.666667
0.384826 0.258819 -0.885958 0.173913 0.583333
0.000000 -1.000000 -0.000000 0.195652 0.000000
0.103114 -0.965926 -0.237392 0.173913 0.083333
0.103114 0.965926 -0.237392 0.173913 0.916667
0.199201 0.866025 -0.458606 0.173913 0.833333
0.345026 -0.500000 -0.794328 0.173913 0.333333
0.281712 -0.707107 -0.648566 0.173913 0.250000
0.398401 -0.000000 -0.917211 0.173913 0.500000
0.281712 0.707107 -0.648566 0.173913 0.750000
0.199201 -0.866025 -0.458606 0.173913 0.166667
0.384826 -0.258819 -0.885958 0.173913 0.416667
0.000000 1.000000 -0.000000 0.195652 1.000000
0.446247 0.707107 -0.548511 0.130435 0.750000
0.546538 0.500000 -0.671786 0.130435 0.666667
0.315544 -0.866025 -0.387856 0.130435 0.166667
0.163338 -0.965926 -0.200769 0.130435 0.083333
0.000000 1.000000 -0.000000 0.152174 1.000000
0.163338 0.965926 -0.200769 0.130435 0.916667
0.609584 -0.258819 -0.749279 0.130435 0.416667
0.546538 -0.500000 -0.671786 0.130435 0.333333
0.609584 0.258819 -0.749279 0.130435 0.583333
0.000000 -1.000000 -0.000000 0.152174 0.000000
0.315544 0.866025 -0.387856 0.130435 0.833333
0.446247 -0.707107 -0.548511 0.130435 0.250000
0.631088 -0.000000 -0.775711 0.130435 0.500000
0.789132 0.258819 -0.557030 0.086957 0.583333
0.816970 -0.000000 -0.576680 0.086957 0.500000
0.408485 0.866025 -0.288340 0.086957 0.833333
0.577685 0.707107 -0.407774 0.086957 0.750000
0.577685 -0.707107 -0.407774 0.086957 0.250000
0.408485 -0.866025 -0.288340 0.086957 0.166667
0.789132 -0.258819 -0.557030 0.086957 0.416667
0.707517 0.500000 -0.499420 0.086957 0.666667
0.211448 -0.965926 -0.149256 0.086957 0.083333
0.000000 1.000000 -0.000000 0.108696 1.000000
0.211447 0.965926 -0.149256 0.086957 0.916667
0.707517 -0.500000 -0.499420 0.086957 0.333333
0.000000 -1.000000 -0.000000 0.108696 0.000000
0.816022 0.500000 -0.290014 0.043478 0.666667
0.910154 0.258819 -0.323469 0.043478 0.583333
0.000000 -1.000000 -0.000000 0.065217 0.000000
0.243875 -0.965926 -0.086673 0.043478 0.083333
0.243875 0.965926 -0.086673 0.043478 0.916667
0.471130 0.866025 -0.167440 0.043478 0.833333
0.816022 -0.500000 -0.290014 0.043478 0.333333
0.666279 -0.707107 -0.236795 0.043478 0.250000
0.942261 -0.000000 -0.334879 0.043478 0.500000
0.666279 0.707107 -0.236795 0.043478 0.750000
0.471131 -0.866025 -0.167440 0.043478 0.166667
0.910154 -0.258819 -0.323469 0.043478 0.416667
0.000000 1.000000 -0.000000 0.065217 1.000000
0.997669 -0.000000 -0.068242 0.000000 0.500000
0.963674 -0.258819 -0.065917 0.000000 0.416667
0.705458 0.707107 -0.048255 0.000000 0.750000
0.864007 0.500000 -0.059099 0.000000 0.666667
0.498834 -0.866025 -0.034121 0.000000 0.166667
0.258216 -0.965926 -0.017662 0.000000 0.083333
0.000000 1.000000 -0.000000 0.021739 1.000000
0.258216 0.965926 -0.017662 0.000000 0.916667
0.864006 -0.500000 -0.059099 0.000000 0.333333
0.963674 0.258819 -0.065917 0.000000 0.583333
0.000000 -1.000000 -0.000000 0.021739 0.000000
0.498834 0.866025 -0.034121 0.000000 0.833333
0.705458 -0.707107 -0.048255 0.000000 0.250000
0.997669 -0.000000 -0.068242 1.000000 0.500000
0.963674 0.258819 -0.065917 1.000000 0.583333
0.945722 0.258819 0.196524 0.956522 0.583333
0.979084 -0.000000 0.203456 0.956522 0.500000
0.705458 0.707107 -0.048255 1.000000 0.750000
0.498834 0.866025 -0.034121 1.000000 0.833333
0.489542 0.866025 0.101728 0.956522 0.833333
0.692317 0.707107 0.143865 0.956522 0.750000
0.498834 -0.866025 -0.034121 1.000000 0.166667
0.705458 -0.707107 -0.048255 1.000000 0.250000
0.692317 -0.707107 0.143865 0.956522 0.250000
0.489542 -0.866025 0.101728 0.956522 0.166667
0.963674 -0.258819 -0.065917 1.000000 0.416667
0.945722 -0.258819 0.196524 0.956522 0.416667
0.864007 0.500000 -0.059099 1.000000 0.666667
0.847912 0.500000 0.176198 0.956522 0.666667
0.258216 -0.965926 -0.017662 1.000000 0.083333
0.253406 -0.965926 0.052658 0.956522 0.083333
0.258216 0.965926 -0.017662 1.000000 0.916667
0.000000 1.000000 -0.000000 0.978261 1.000000
0.253406 0.965926 0.052658 0.956522 0.916667
0.864006 -0.500000 -0.059099 1.000000 0.333333
0.847911 -0.500000 0.176198 0.956522 0.333333
0.000000 -1.000000 -0.000000 0.978261 0.000000
0.857631 -0.258819 0.444389 0.913043 0.416667
0.768931 -0.500000 0.398428 0.913043 0.333333
0.768931 0.500000 0.398428 0.913043 0.666667
0.857631 0.258819 0.444389 0.913043 0.583333
0.000000 -1.000000 -0.000000 0.934783 0.000000
0.229802 -0.965926 0.119074 0.913043 0.083333
0.229802 0.965926 0.119074 0.913043 0.916667
0.443942 0.866025 0.230033 0.913043 0.833333
0.627830 -0.707107 0.325315 0.913043 0.250000
0.887885 -0.000000 0.460065 0.913043 0.500000
0.627830 0.707107 0.325315 0.913043 0.750000
0.443943 -0.866025 0.230033 0.913043 0.166667
0.000000 1.000000 -0.000000 0.934783 1.000000
0.730836 -0.000000 0.682553 0.869565 0.500000
0.705933 -0.258819 0.659296 0.869565 0.416667
0.516779 0.707107 0.482638 0.869565 0.750000
0.632922 0.500000 0.591109 0.869565 0.666667
0.365418 -0.866025 0.341277 0.869565 0.166667
0.189154 -0.965926 0.176658 0.869565 0.083333
0.000000 1.000000 -0.000000 0.891304 1.000000
0.189154 0.965926 0.176658 0.869565 0.916667
0.632922 -0.500000 0.591108 0.869565 0.333333
0.705933 0.258819 0.659296 0.869565 0.583333
0.000000 -1.000000 -0.000000 0.891304 0.000000
0.365418 0.866025 0.341277 0.869565 0.833333
0.516779 -0.707107 0.482638 0.869565 0.250000
0.449973 -0.500000 0.739949 0.826087 0.333333
0.367401 -0.707107 0.604166 0.826087 0.250000
0.501879 0.258819 0.825306 0.826087 0.583333
0.519583 -0.000000 0.854420 0.826087 0.500000
0.259792 0.866025 0.427210 0.826087 0.833333
0.367401 0.707107 0.604166 0.826087 0.750000
0.259792 -0.866025 0.427210 0.826087 0.166667
0.501879 -0.258819 0.825306 0.826087 0.416667
0.449973 0.500000 0.739949 0.826087 0.666667
0.134478 -0.965926 0.221140 0.826087 0.083333
0.000000 1.000000 -0.000000 0.847826 1.000000
0.134478 0.965926 0.221140 0.826087 0.916667
0.000000 -1.000000 -0.000000 0.847826 0.000000
0.000000 1.000000 -0.000000 0.804348 1.000000
0.069828 0.965926 0.249221 0.782609 0.916667
0.260603 -0.258819 0.930106 0.782609 0.416667
0.233650 -0.500000 0.833911 0.782609 0.333333
0.233650 0.500000 0.833911 0.782609 0.666667
0.260603 0.258819 0.930106 0.782609 0.583333
0.000000 -1.000000 -0.000000 0.804348 0.000000
0.069828 -0.965926 0.249221 0.782609 0.083333
0.134898 0.866025 0.481459 0.782609 0.833333
0.190775 -0.707107 0.680885 0.782609 0.250000
0.269796 -0.000000 0.962917 0.782609 0.500000
0.190775 0.707107 0.680885 0.782609 0.750000
0.134898 -0.866025 0.481459 0.782609 0.166667
0.000000 1.000000 -0.000000 0.760870 1.000000
0.000000 -1.000000 -0.000000 0.760870 0.000000
4 0 1 2 3
4 4 5 6 7
4 8 9 10 11
4 12 0 3 13
4 14 4 7 15
3 16 17 18
4 19 8 11 20
4 21 12 13 22
3 23 14 15
4 1 16 18 2
4 5 19 20 6
4 9 21 22 10
4 20 11 24 25
4 22 13 26 27
3 28 15 29
4 2 18 30 31
4 6 20 25 32
4 10 22 27 33
4 3 2 31 34
4 7 6 32 35
4 11 10 33 24
4 13 3 34 26
4 15 7 35 29
3 18 36 30
4 24 33 37 38
4 26 34 39 40
4 29 35 41 42
3 30 43 44
4 25 24 38 45
4 27 26 40 46
3 47 29 42
4 31 30 44 48
4 32 25 45 49
4 33 27 46 37
4 34 31 48 39
4 35 32 49 41
4 49 45 50 51
4 37 46 52 53
4 39 48 54 55
4 41 49 51 56
4 38 37 53 57
4 40 39 55 58
4 42 41 56 59
3 44 60 61
4 45 38 57 50
4 46 40 58 52
3 62 42 59
4 48 44 61 54
3 61 63 64
4 50 57 65 66
4 52 58 67 68
3 69 59 70
4 54 61 64 71
4 51 50 66 72
4 53 52 68 73
4 55 54 71 74
4 56 51 72 75
4 57 53 73 65
4 58 55 74 67
4 59 56 75 70
4 75 72 76 77
4 65 73 78 79
4 67 74 80 81
4 70 75 77 82
3 64 83 84
4 66 65 79 85
4 68 67 81 86
3 87 70 82
4 71 64 84 88
4 72 66 85 76
4 73 68 86 78
4 74 71 88 80
4 88 84 89 90
4 76 85 91 92
4 78 86 93 94
4 80 88 90 95
4 77 76 92 96
4 79 78 94 97
4 81 80 95 98
4 82 77 96 99
3 84 100 89
4 85 79 97 91
4 86 81 98 93
3 101 82 99
4 99 96 102 103
3 89 104 105
4 91 97 106 107
4 93 98 108 109
3 110 99 103
4 90 89 105 111
4 92 91 107 112
4 94 93 109 113
4 95 90 111 114
4 96 92 112 102
4 97 94 113 106
4 98 95 114 108
4 114 111 115 116
4 102 112 117 118
4 106 113 119 120
4 108 114 116 121
4 103 102 118 122
3 105 123 124
4 107 106 120 125
4 109 108 121 126
3 127 103 122
4 111 105 124 115
4 112 107 125 117
4 113 109 126 119
3 128 122 129
4 115 124 130 131
4 117 125 132 133
4 119 126 134 135
4 116 115 131 136
4 118 117 133 137
4 120 119 135 138
4 121 116 136 139
4 122 118 137 129
3 124 140 130
4 125 120 138 132
4 126 121 139 134
4 139 136 141 142
4 129 137 143 144
3 130 145 146
4 132 138 147 148
4 134 139 142 149
3 150 129 144
4 131 130 146 151
4 133 132 148 152
4 135 134 149 153
4 136 131 151 141
4 137 133 152 143
4 138 135 153 147
4 141 151 154 155
4 143 152 156 157
4 147 153 158 159
4 142 141 155 160
4 144 143 157 161
3 146 162 163
4 148 147 159 164
4 149 142 160 165
3 166 144 161
4 151 146 163 154
4 152 148 164 156
4 153 149 165 158
4 165 160 167 168
3 169 161 170
4 154 163 171 172
4 156 164 173 174
4 158 165 168 175
4 155 154 172 176
4 157 156 174 177
4 159 158 175 178
4 160 155 176 167
4 161 157 177 170
3 163 179 171
4 164 159 178 173
4 167 176 180 181
4 170 177 182 183
3 171 184 185
4 173 178 186 187
4 168 167 181 188
3 189 170 183
4 172 171 185 190
4 174 173 187 191
4 175 168 188 192
4 176 172 190 180
4 177 174 191 182
4 178 175 192 186
4 192 188 193 194
4 180 190 195 196
4 182 191 197 198
4 186 192 194 199
4 181 180 196 200
4 183 182 198 201
3 185 202 203
4 187 186 199 204
4 188 181 200 193
3 205 183 201
4 190 185 203 195
4 191 187 204 197
4 193 200 206 207
3 208 201 209
4 195 203 210 211
4 197 204 212 213
4 194 193 207 214
4 196 195 211 215
4 198 197 213 216
4 199 194 214 217
4 200 196 215 206
4 201 198 216 209
3 203 218 210
4 204 199 217 212
4 217 214 219 220
4 206 215 221 222
4 209 216 223 224
3 210 225 226
4 212 217 220 227
4 207 206 222 228
3 229 209 224
4 211 210 226 230
4 213 212 227 231
4 214 207 228 219
4 215 211 230 221
4 216 213 231 223
4 232 233 234 235
4 236 237 238 239
4 240 241 242 243
4 244 232 235 245
4 246 236 239 247
4 248 240 243 249
3 250 251 252
4 253 244 245 254
4 233 246 247 234
3 255 248 249
4 237 250 252 238
4 241 253 254 242
4 254 245 256 257
4 234 247 258 259
3 260 249 261
4 238 252 262 263
4 242 254 257 264
4 235 234 259 265
4 239 238 263 266
4 243 242 264 267
4 245 235 265 256
4 247 239 266 258
4 249 243 267 261
3 252 268 262
4 256 265 269 270
4 258 266 271 272
4 261 267 273 274
3 262 275 276
4 257 256 270 277
4 259 258 272 278
3 279 261 274
4 263 262 276 280
4 264 257 277 281
4 265 259 278 269
4 266 263 280 271
4 267 264 281 273
4 281 277 282 283
4 269 278 284 285
4 271 280 286 287
4 273 281 283 288
4 270 269 285 289
4 272 271 287 290
4 274 273 288 291
3 276 292 293
4 277 270 289 282
4 278 272 290 284
3 294 274 291
4 280 276 293 286
3 293 295 296
4 282 289 297 298
4 284 290 299 300
3 301 291 302
4 286 293 296 303
4 283 282 298 304
4 285 284 300 305
4 287 286 303 306
4 288 283 304 307
4 289 285 305 297
4 290 287 306 299
4 291 288 307 302
4 307 304 5 4
4 297 305 9 8
4 299 306 0 12
4 302 307 4 14
3 296 308 16
4 298 297 8 19
4 300 299 12 21
3 309 302 14
4 303 296 16 1
4 304 298 19 5
4 305 300 21 9
4 306 303 1 0

101
shaders/normal.comp
Unescape View File

@ -3,85 +3,88 @@
layout (local_size_x = 32) in; layout (local_size_x = 32) in;
struct Vertex { struct Vertex {
vec3 position; vec3 position;
vec3 color; vec3 color;
uvec3 normal; uvec3 normal;
vec3 velocity;
vec3 prevPosition;
float invM;
}; };
struct Face { struct Face {
uint a; uint a;
uint b; uint b;
uint c; uint c;
}; };
layout (std430, set = 0, binding = 0) buffer VertexBuffer { layout (std430, set = 0, binding = 0) buffer VertexBuffer {
Vertex vertices[]; Vertex vertices[];
}; };
layout (std430, set = 0, binding = 1) buffer FaceBuffer { layout (std430, set = 0, binding = 1) buffer FaceBuffer {
Face faces[]; Face faces[];
}; };
layout (std140, set = 0, binding = 5) uniform Sizes { layout (std140, set = 0, binding = 5) uniform Sizes {
uint vertexCount; uint vertexCount;
uint faceCount; uint faceCount;
}; };
layout (push_constant) uniform PushConstants { layout (push_constant) uniform PushConstants {
uint state; uint state;
}; };
void atomicAddVec3(uint vID, vec3 add){ void atomicAddVec3(uint vID, vec3 add){
for (int i = 0; i < 3; i++) { for (int i = 0; i < 3; i++) {
uint expected_memory = vertices[vID].normal[i]; uint expected_memory = vertices[vID].normal[i];
float floatInput = uintBitsToFloat(vertices[vID].normal[i]) + add[i]; float floatInput = uintBitsToFloat(vertices[vID].normal[i]) + add[i];
uint actual_content = atomicCompSwap(vertices[vID].normal[i], expected_memory, floatBitsToUint(floatInput)); uint actual_content = atomicCompSwap(vertices[vID].normal[i], expected_memory, floatBitsToUint(floatInput));
while (actual_content != expected_memory) { while (actual_content != expected_memory) {
expected_memory = actual_content; expected_memory = actual_content;
floatInput = uintBitsToFloat(expected_memory) + add[i]; floatInput = uintBitsToFloat(expected_memory) + add[i];
actual_content = atomicCompSwap(vertices[vID].normal[i], expected_memory, floatBitsToUint(floatInput)); actual_content = atomicCompSwap(vertices[vID].normal[i], expected_memory, floatBitsToUint(floatInput));
} }
} }
} }
void reset(uint vID){ void reset(uint vID){
vertices[vID].normal = floatBitsToUint(vec3(0, 0, 0)); vertices[vID].normal = floatBitsToUint(vec3(0, 0, 0));
} }
void accumulate(uint fID){ void accumulate(uint fID){
Face f = faces[fID]; Face f = faces[fID];
Vertex v1 = vertices[f.a]; Vertex v1 = vertices[f.a];
Vertex v2 = vertices[f.b]; Vertex v2 = vertices[f.b];
Vertex v3 = vertices[f.c]; Vertex v3 = vertices[f.c];
vec3 weightedNormal = cross(v3.position - v1.position, v2.position - v1.position); vec3 weightedNormal = cross(v3.position - v1.position, v2.position - v1.position);
atomicAddVec3(f.a, weightedNormal); atomicAddVec3(f.a, weightedNormal);
atomicAddVec3(f.b, weightedNormal); atomicAddVec3(f.b, weightedNormal);
atomicAddVec3(f.c, weightedNormal); atomicAddVec3(f.c, weightedNormal);
} }
void norm(uint vID){ void norm(uint vID){
vertices[vID].normal = floatBitsToUint(normalize(uintBitsToFloat(vertices[vID].normal))); vertices[vID].normal = floatBitsToUint(normalize(uintBitsToFloat(vertices[vID].normal)));
} }
void main() { void main() {
uint id = gl_GlobalInvocationID.x; uint id = gl_GlobalInvocationID.x;
switch (state){ switch (state){
case 0: case 0:
if (id < vertexCount){ if (id < vertexCount){
reset(id); reset(id);
} }
break; break;
case 1: case 1:
if (id < faceCount){ if (id < faceCount){
accumulate(id); accumulate(id);
} }
break; break;
case 2: case 2:
if (id < vertexCount){ if (id < vertexCount){
norm(id); norm(id);
} }
break; break;
} }
} }

109
shaders/pbd.comp
Unescape View File

@ -3,61 +3,110 @@
layout (local_size_x = 32) in; layout (local_size_x = 32) in;
struct Vertex { struct Vertex {
vec3 position; vec3 position;
vec3 color; vec3 color;
vec3 normal; vec3 normal;
vec3 velocity;
vec3 prevPosition;
float w;
}; };
struct Edge { struct Edge {
uint a; uint a;
uint b; uint b;
float restLength;
}; };
layout (std430, set = 0, binding = 0) buffer VertexBuffer { layout (std430, set = 0, binding = 0) buffer VertexBuffer {
Vertex vertices[]; Vertex vertices[];
}; };
layout (std430, set = 0, binding = 2) buffer EdgeBuffer { layout (std430, set = 0, binding = 2) buffer EdgeBuffer {
Edge edges[]; Edge edges[];
}; };
layout (std140, set = 0, binding = 5) uniform Sizes { layout (std140, set = 0, binding = 5) uniform Sizes {
uint vertexCount; uint vertexCount;
uint faceCount; uint faceCount;
};
layout (std140, set = 1, binding = 0) uniform Properties {
vec3 gravity;
float dt;
uint k;
}; };
layout (std140, set = 1, binding = 0) uniform UBO { struct Partition {
float dt; uint offset;
vec3 gravity; uint size;
}; };
layout (push_constant) uniform PushConstants { layout (push_constant, std430) uniform PushConstants {
uint state; uint state;
int vertexOffset; Partition edgePartition;
Partition tetrahedronPartition;
}; };
void preSolve(uint vID){ void preSolve(uint vID){
vertices[vID].position += vec3(0.0001, 0, 0); vertices[vID].prevPosition = vertices[vID].position;
// vertices[vID].velocity += dt * gravity;
vertices[vID].position += dt * vertices[vID].velocity;
}
void solveEdge(uint eID){
Edge e = edges[eID];
Vertex v1 = vertices[e.a];
Vertex v2 = vertices[e.b];
vec3 diff = v1.position - v2.position;
float currentLength = length(diff);
float alpha = 0.0001 / (dt * dt);
float s = (currentLength - e.restLength) / (1 + 1 + alpha);
vec3 g1 = normalize(diff);
vec3 g2 = -g1;
vec3 delta1 = -s * 1 * g1;
vec3 delta2 = -s * 1 * g2;
vertices[e.a].position += delta1;
vertices[e.b].position += delta2;
}
void solveTetrahedron(uint tetID){
} }
void postSolve(uint vID){ void postSolve(uint vID){
vertices[vID].position += vec3(0, 0.0001, 0); vertices[vID].velocity = (vertices[vID].position - vertices[vID].prevPosition) / dt;
} }
void main() { void main() {
uint id = gl_GlobalInvocationID.x; uint id = gl_GlobalInvocationID.x;
switch (state){ switch (state){
case 0: case 0:
if (id < vertexCount){ if (id < vertexCount){
preSolve(id); preSolve(id);
} }
break; break;
case 2: case 1:
if (id < vertexCount){ if (id < edgePartition.size){
postSolve(id); solveEdge(id + edgePartition.offset);
} }
break; if (id < tetrahedronPartition.size){
} solveTetrahedron(id + tetrahedronPartition.offset);
}
break;
case 2:
if (id < vertexCount){
postSolve(id);
}
break;
}
} }

22
src/constraints.cpp
Unescape View File

@ -1 +1,23 @@
#include "constraints.hpp" #include "constraints.hpp"
void ConstraintData::recordNewPartition() {
prePartitionEdgeCount = edges.size();
prePartitionTetrahedronCount = tetrahedra.size();
}
void ConstraintData::writePartitionInformation() {
edgePartitions.emplace_back(prePartitionEdgeCount, edges.size() - prePartitionEdgeCount);
tetrahedronPartitions.emplace_back(prePartitionTetrahedronCount, tetrahedra.size() - prePartitionTetrahedronCount);
}
void DistanceConstraint::writeData(ConstraintData &dataLists) const {
dataLists.edges.push_back(Edge(a, b, length));
}
void AreaConstraint::writeData(ConstraintData &dataLists) const {
dataLists.triangles.push_back(Triangle(Face(a, b, c), area));
}
void VolumeConstraint::writeData(ConstraintData &dataLists) const {
dataLists.tetrahedra.push_back(Tetrahedron(a, b, c, d, volume));
}

152
src/simulation.cpp
Unescape View File

@ -12,14 +12,6 @@
Simulation::Simulation() { Simulation::Simulation() {
createMeshBuffers(); createMeshBuffers();
struct SizeInformation {
uint32_t vertexCount;
uint32_t faceCount;
uint32_t edgeCount;
uint32_t triangleCount;
uint32_t tetrahedronCount;
};
SizeInformation sizeInformation {}; SizeInformation sizeInformation {};
sizeInformation.vertexCount = vertexBuffer->size / sizeof(Vertex); sizeInformation.vertexCount = vertexBuffer->size / sizeof(Vertex);
sizeInformation.faceCount = faceBuffer->size / sizeof(Face); sizeInformation.faceCount = faceBuffer->size / sizeof(Face);
@ -29,6 +21,15 @@ Simulation::Simulation() {
VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT, VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT,
VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE, 0); VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE, 0);
properties.gravity = {0, -9.81, 0};
properties.k = 10;
properties.dt = 1.f / 144.f / static_cast<float>(properties.k);
propertiesBuffer = make_unique<Buffer>(
sizeof(Properties), &properties, sizeof(properties),
VK_BUFFER_USAGE_UNIFORM_BUFFER_BIT,
VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE, 0);
descriptorPool->bindBuffer(*vertexBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 0); descriptorPool->bindBuffer(*vertexBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 0);
descriptorPool->bindBuffer(*faceBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 1); descriptorPool->bindBuffer(*faceBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 1);
descriptorPool->bindBuffer(*edgeBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 2); descriptorPool->bindBuffer(*edgeBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 2);
@ -36,6 +37,8 @@ Simulation::Simulation() {
descriptorPool->bindBuffer(*tetrahedronBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 4); descriptorPool->bindBuffer(*tetrahedronBuffer, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, DescriptorSet::MESH, 4);
descriptorPool->bindBuffer(*sizeInformationBuffer, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, DescriptorSet::MESH, 5); descriptorPool->bindBuffer(*sizeInformationBuffer, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, DescriptorSet::MESH, 5);
descriptorPool->bindBuffer(*propertiesBuffer, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, DescriptorSet::SIMULATION, 0);
createComputePipelines(); createComputePipelines();
char* stats; char* stats;
@ -45,8 +48,8 @@ Simulation::Simulation() {
} }
void Simulation::createMeshBuffers() { void Simulation::createMeshBuffers() {
Mesh sphere("models/sphere_high.ply"); Mesh sphere("models/sphere_medium.ply");
Mesh bunny("models/bunny_high.ply"); Mesh bunny("models/bunny_medium.ply");
auto body = std::make_unique<SoftBody>(&sphere, 0.3f); auto body = std::make_unique<SoftBody>(&sphere, 0.3f);
@ -64,40 +67,79 @@ void Simulation::createMeshBuffers() {
} }
vector<Vertex> vertices; vector<Vertex> vertices;
vector<Edge> edges;
vector<Triangle> triangles;
vector<Face> faces; vector<Face> faces;
vector<Tetrahedron> tetrahedra;
for (std::unique_ptr<SoftBody> &softBody : softBodies){ for (std::unique_ptr<SoftBody> &currentSoftBody : softBodies){
softBody->firstIndex = faces.size() * 3; currentSoftBody->firstIndex = faces.size() * 3;
softBody->vertexOffset = static_cast<int32_t>(vertices.size()); currentSoftBody->vertexOffset = static_cast<int32_t>(vertices.size());
int32_t vertexOffset = currentSoftBody->vertexOffset;
int32_t off = softBody->vertexOffset; for (auto &face : currentSoftBody->faces){
face.a += vertexOffset;
face.b += vertexOffset;
face.c += vertexOffset;
}
for (auto &edge : softBody->edges){ for (auto &edge : currentSoftBody->constraintData.edges){
edge.a += off; edge.a += vertexOffset;
edge.b += off; edge.b += vertexOffset;
} }
for (auto &face : softBody->faces){ for (auto &triangle : currentSoftBody->constraintData.triangles){
face.a += off; triangle.a += vertexOffset;
face.b += off; triangle.b += vertexOffset;
face.c += off; triangle.c += vertexOffset;
} }
for (auto &tetrahedron : softBody->tetrahedra){ for (auto &tetrahedron : currentSoftBody->constraintData.tetrahedra){
tetrahedron.a += off; tetrahedron.a += vertexOffset;
tetrahedron.b += off; tetrahedron.b += vertexOffset;
tetrahedron.c += off; tetrahedron.c += vertexOffset;
tetrahedron.d += off; tetrahedron.d += vertexOffset;
} }
vertices.insert(vertices.end(), softBody->vertices.begin(), softBody->vertices.end()); vertices.insert(vertices.end(), currentSoftBody->vertices.begin(), currentSoftBody->vertices.end());
edges.insert(edges.end(), softBody->edges.begin(), softBody->edges.end()); faces.insert(faces.end(), currentSoftBody->faces.begin(), currentSoftBody->faces.end());
triangles.insert(triangles.end(), softBody->triangles.begin(), softBody->triangles.end());
faces.insert(faces.end(), softBody->faces.begin(), softBody->faces.end()); constraintData.partitionCount = std::max(constraintData.partitionCount, currentSoftBody->constraintData.partitionCount);
tetrahedra.insert(tetrahedra.end(), softBody->tetrahedra.begin(), softBody->tetrahedra.end());
auto combine = [&currentSoftBody, this] (
auto &globalIndices, auto &bodyIndices,
vector<ConstraintData::Partition> &globalPartitions, vector<ConstraintData::Partition> &bodyPartitions){
if (globalPartitions.size() < currentSoftBody->constraintData.partitionCount)
globalPartitions.resize(currentSoftBody->constraintData.partitionCount);
uint32_t offsetAdded = 0;
for (uint32_t partition = 0; partition < constraintData.partitionCount; partition++){
ConstraintData::Partition &globalPartition = globalPartitions[partition];
globalPartition.offset += offsetAdded;
if (partition < bodyPartitions.size()){
const ConstraintData::Partition &bodyPartition = bodyPartitions[partition];
auto dst = globalIndices.begin() + globalPartition.offset;
auto srcStart = bodyIndices.begin() + bodyPartition.offset;
uint32_t count = bodyPartition.size;
globalIndices.insert(dst, srcStart, srcStart + count);
globalPartition.size += count;
offsetAdded += count;
}
}
};
combine(constraintData.edges, currentSoftBody->constraintData.edges,
constraintData.edgePartitions, currentSoftBody->constraintData.edgePartitions);
combine(constraintData.tetrahedra, currentSoftBody->constraintData.tetrahedra,
constraintData.tetrahedronPartitions, currentSoftBody->constraintData.tetrahedronPartitions);
constraintData.triangles.insert(constraintData.triangles.end(), currentSoftBody->constraintData.triangles.begin(), currentSoftBody->constraintData.triangles.end());
constraintData.tetrahedra.insert(constraintData.tetrahedra.end(), currentSoftBody->constraintData.tetrahedra.begin(), currentSoftBody->constraintData.tetrahedra.end());
} }
class SimulationBuffer : public Buffer { class SimulationBuffer : public Buffer {
@ -106,11 +148,13 @@ void Simulation::createMeshBuffers() {
: Buffer(size, data, size, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | additionalUsageFlags, VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE, 0) {} : Buffer(size, data, size, VK_BUFFER_USAGE_STORAGE_BUFFER_BIT | additionalUsageFlags, VMA_MEMORY_USAGE_AUTO_PREFER_DEVICE, 0) {}
}; };
vertices[0].position += glm::vec3(0, 0.2, 0);
vertexBuffer = make_unique<SimulationBuffer>(vertices.data(), vertices.size() * sizeof(Vertex), VK_BUFFER_USAGE_VERTEX_BUFFER_BIT); vertexBuffer = make_unique<SimulationBuffer>(vertices.data(), vertices.size() * sizeof(Vertex), VK_BUFFER_USAGE_VERTEX_BUFFER_BIT);
faceBuffer = make_unique<SimulationBuffer>(faces.data(), faces.size() * sizeof(Face), VK_BUFFER_USAGE_INDEX_BUFFER_BIT); faceBuffer = make_unique<SimulationBuffer>(faces.data(), faces.size() * sizeof(Face), VK_BUFFER_USAGE_INDEX_BUFFER_BIT);
edgeBuffer = make_unique<SimulationBuffer>(edges.data(), edges.size() * sizeof(Edge)); edgeBuffer = make_unique<SimulationBuffer>(constraintData.edges.data(), constraintData.edges.size() * sizeof(Edge));
triangleBuffer = make_unique<SimulationBuffer>(triangles.data(), triangles.size() * sizeof(Triangle)); triangleBuffer = make_unique<SimulationBuffer>(constraintData.triangles.data(), constraintData.triangles.size() * sizeof(Triangle));
tetrahedronBuffer = make_unique<SimulationBuffer>(tetrahedra.data(), tetrahedra.size() * sizeof(Tetrahedron)); tetrahedronBuffer = make_unique<SimulationBuffer>(constraintData.tetrahedra.data(), constraintData.tetrahedra.size() * sizeof(Tetrahedron));
} }
void Simulation::recordDrawCommands() { void Simulation::recordDrawCommands() {
@ -144,25 +188,38 @@ void Simulation::recordComputeCommands(VkCommandBuffer cmdBuffer) {
vkCmdBindPipeline(cmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, pbdPipeline->handle); vkCmdBindPipeline(cmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, pbdPipeline->handle);
vkCmdBindDescriptorSets(cmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, pbdPipeline->layout, 0, 1, &descriptorPool->sets[DescriptorSet::MESH], 0, nullptr); vkCmdBindDescriptorSets(cmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, pbdPipeline->layout, 0, 1, &descriptorPool->sets[DescriptorSet::MESH], 0, nullptr);
vkCmdBindDescriptorSets(cmdBuffer, VK_PIPELINE_BIND_POINT_COMPUTE, pbdPipeline->layout, 1, 1, &descriptorPool->sets[DescriptorSet::SIMULATION], 0, nullptr);
uint32_t state; uint32_t state;
size_t subSteps = 10; for (size_t i = 0; i < properties.k; i++){
for (size_t i = 0; i < subSteps; i++){
state = 0; state = 0;
vkCmdPushConstants(cmdBuffer, pbdPipeline->layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &state); vkCmdPushConstants(cmdBuffer, pbdPipeline->layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &state);
vkCmdDispatch(cmdBuffer, vertexGroupCount, 1, 1); vkCmdDispatch(cmdBuffer, vertexGroupCount, 1, 1);
vkCmdPipelineBarrier(cmdBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &barrier, 0, nullptr, 0, nullptr); vkCmdPipelineBarrier(cmdBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &barrier, 0, nullptr, 0, nullptr);
/* state = 1;
uint32_t partitionCount = 1; vkCmdPushConstants(cmdBuffer, pbdPipeline->layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &state);
for (uint32_t partition = 0; partition < partitionCount; partition++){
uint32_t partitionSize = 1; for (uint32_t partition = 0; partition < constraintData.partitionCount; partition++){
vkCmdDispatch(cmdBuffer, partitionSize, 1, 1); auto edgePartition = constraintData.edgePartitions[partition];
auto tetrahedronPartition = constraintData.tetrahedronPartitions[partition];
if (edgePartition.size == 0 && tetrahedronPartition.size == 0)
continue;
ConstraintData::Partition partitions[2] = {edgePartition, tetrahedronPartition};
vkCmdPushConstants(cmdBuffer, pbdPipeline->layout, VK_SHADER_STAGE_COMPUTE_BIT,
offsetof(PBDPushData, edgePartition),
sizeof(partitions), partitions);
uint32_t invocations = getGroupCount(std::max(edgePartition.size, tetrahedronPartition.size), BlOCK_SIZE);
vkCmdDispatch(cmdBuffer, invocations, 1, 1);
vkCmdPipelineBarrier(cmdBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &barrier, 0, nullptr, 0, nullptr); vkCmdPipelineBarrier(cmdBuffer, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, 1, &barrier, 0, nullptr, 0, nullptr);
} }
*/
state = 2; state = 2;
vkCmdPushConstants(cmdBuffer, pbdPipeline->layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &state); vkCmdPushConstants(cmdBuffer, pbdPipeline->layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(uint32_t), &state);
vkCmdDispatch(cmdBuffer, vertexGroupCount, 1, 1); vkCmdDispatch(cmdBuffer, vertexGroupCount, 1, 1);
@ -195,10 +252,11 @@ void Simulation::createComputePipelines() {
{ {
layouts.push_back(descriptorPool->layouts[DescriptorSet::MESH]); layouts.push_back(descriptorPool->layouts[DescriptorSet::MESH]);
layouts.push_back(descriptorPool->layouts[DescriptorSet::SIMULATION]);
pushRanges.push_back({ pushRanges.push_back({
.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT, .stageFlags = VK_SHADER_STAGE_COMPUTE_BIT,
.offset = 0, .offset = 0,
.size = 2 * sizeof(uint32_t) .size = sizeof(PBDPushData)
}); });
pbdPipeline = unique_ptr<ComputePipeline>(new ComputePipeline("shaders/pbd.spv", layouts, pushRanges)); pbdPipeline = unique_ptr<ComputePipeline>(new ComputePipeline("shaders/pbd.spv", layouts, pushRanges));
} }

130
src/soft_body.cpp
Unescape View File

@ -3,10 +3,13 @@
#include "mesh.hpp" #include "mesh.hpp"
#include "constraints.hpp" #include "constraints.hpp"
#include "tetgen.h" #include "tetgen.h"
#include "timer.hpp"
#include <omp.h>
#include <iostream>
SoftBody::SoftBody(Mesh* mesh, float compliance) : compliance(compliance) { SoftBody::SoftBody(Mesh* mesh, float compliance) : compliance(compliance) {
tetgenbehavior behavior; tetgenbehavior behavior;
behavior.parse_commandline(std::string("pYa0.01Qfez").data()); behavior.parse_commandline(std::string("pYa0.001Qfez").data());
tetgenio in; tetgenio in;
in.numberofpoints = static_cast<int>(mesh->vertices.size()); in.numberofpoints = static_cast<int>(mesh->vertices.size());
@ -57,34 +60,35 @@ SoftBody::SoftBody(Mesh* mesh, float compliance) : compliance(compliance) {
} }
faces.reserve(out.numberoftrifaces); faces.reserve(out.numberoftrifaces);
triangles.reserve(out.numberoftrifaces); constraintData.triangles.reserve(out.numberoftrifaces);
for (size_t i = 0; i < out.numberoftrifaces; i++){ for (size_t i = 0; i < out.numberoftrifaces; i++){
uint32_t a = out.trifacelist[i * 3 + 0]; uint32_t a = out.trifacelist[i * 3 + 0];
uint32_t b = out.trifacelist[i * 3 + 1]; uint32_t b = out.trifacelist[i * 3 + 1];
uint32_t c = out.trifacelist[i * 3 + 2]; uint32_t c = out.trifacelist[i * 3 + 2];
if (out.trifacemarkerlist[i] != 0) if (out.trifacemarkerlist[i] != 0)
faces.emplace_back(Face(a, b, c)); faces.emplace_back(Face(a, b, c));
triangles.emplace_back(Triangle(Face(a, b, c))); constraintData.triangles.emplace_back(Triangle(Face(a, b, c), 0));
} }
faces.shrink_to_fit(); faces.shrink_to_fit();
edges.reserve(out.numberofedges); constraintData.edges.reserve(out.numberofedges);
for (size_t i = 0; i < out.numberofedges; i++) { for (size_t i = 0; i < out.numberofedges; i++) {
uint32_t a = out.edgelist[i * 2 + 0]; uint32_t a = out.edgelist[i * 2 + 0];
uint32_t b = out.edgelist[i * 2 + 1]; uint32_t b = out.edgelist[i * 2 + 1];
float length = glm::length(vertices[a].position - vertices[b].position); float length = glm::length(vertices[a].position - vertices[b].position);
edges.emplace_back(Edge(a, b, length)); constraintData.edges.emplace_back(Edge(a, b, length));
} }
tetrahedra.reserve(out.numberoftetrahedra); constraintData.tetrahedra.reserve(out.numberoftetrahedra);
for (size_t i = 0; i < out.numberoftetrahedra; i++){ for (size_t i = 0; i < out.numberoftetrahedra; i++){
uint32_t a = out.tetrahedronlist[i * 4 + 0]; uint32_t a = out.tetrahedronlist[i * 4 + 0];
uint32_t b = out.tetrahedronlist[i * 4 + 1]; uint32_t b = out.tetrahedronlist[i * 4 + 1];
uint32_t c = out.tetrahedronlist[i * 4 + 2]; uint32_t c = out.tetrahedronlist[i * 4 + 2];
uint32_t d = out.tetrahedronlist[i * 4 + 3]; uint32_t d = out.tetrahedronlist[i * 4 + 3];
tetrahedra.emplace_back(Tetrahedron(a, b, c, d)); constraintData.tetrahedra.emplace_back(Tetrahedron(a, b, c, d, 0));
} }
splitConstraints();
} }
void SoftBody::applyVertexOffset(const glm::vec3 &offset) { void SoftBody::applyVertexOffset(const glm::vec3 &offset) {
@ -93,6 +97,116 @@ void SoftBody::applyVertexOffset(const glm::vec3 &offset) {
} }
} }
void SoftBody::applyIndexOffset(int32_t offset) { vector<unordered_set<const Constraint *>> createPartitions(const Graph &graph){
unordered_map<const Constraint *, uint32_t> constraintToColor;
vector<unordered_set<const Constraint *>> colorToConstraintList;
for (const auto& [constraint, adjacentList] : graph){
unordered_set<uint32_t> forbiddenColors;
for (auto adjacent : adjacentList)
if (constraintToColor.contains(adjacent))
forbiddenColors.insert(constraintToColor.at(adjacent));
uint32_t assignColor;
if (forbiddenColors.size() == colorToConstraintList.size()) {
assignColor = forbiddenColors.size();
colorToConstraintList.emplace_back();
} else {
uint32_t minColor;
uint32_t minColorUseCount = UINT32_MAX;
for (uint32_t color = 0; color < colorToConstraintList.size(); color++){
if (colorToConstraintList[color].size() < minColorUseCount && !forbiddenColors.contains(color)){
minColorUseCount = colorToConstraintList[color].size();
minColor = color;
}
}
assignColor = minColor;
}
constraintToColor[constraint] = assignColor;
colorToConstraintList[assignColor].insert(constraint);
}
return colorToConstraintList;
}
void SoftBody::splitConstraints() {
omp_set_num_threads(omp_get_num_procs() - 2);
unordered_map<uint32_t, vector<const Constraint *>> pointToConstraints;
vector<DistanceConstraint> lengthConstraints;
vector<VolumeConstraint> volumeConstraints;
lengthConstraints.reserve(constraintData.edges.size());
volumeConstraints.reserve(constraintData.tetrahedra.size());
for (const Edge &edge : constraintData.edges)
lengthConstraints.push_back(DistanceConstraint(edge));
for (const Tetrahedron &tetrahedron : constraintData.tetrahedra)
volumeConstraints.push_back(VolumeConstraint(tetrahedron));
for (const DistanceConstraint &distanceConstraint : lengthConstraints){
pointToConstraints[distanceConstraint.a].push_back(&distanceConstraint);
pointToConstraints[distanceConstraint.b].push_back(&distanceConstraint);
}
for (const VolumeConstraint &volumeConstraint : volumeConstraints){
pointToConstraints[volumeConstraint.a].push_back(&volumeConstraint);
pointToConstraints[volumeConstraint.b].push_back(&volumeConstraint);
pointToConstraints[volumeConstraint.c].push_back(&volumeConstraint);
pointToConstraints[volumeConstraint.d].push_back(&volumeConstraint);
}
Graph graph;
auto setNeighbors = [&pointToConstraints, &graph](const Constraint * constraint, const vector<uint32_t> &vIDs){
unordered_set<const Constraint *> neighbors;
for (uint32_t vID : vIDs)
neighbors.insert(pointToConstraints[vID].begin(), pointToConstraints[vID].end());
neighbors.erase(constraint);
#pragma omp critical
graph[constraint].insert(graph[constraint].end(), neighbors.begin(), neighbors.end());
};
#pragma omp parallel for default(none) shared(setNeighbors, lengthConstraints)
for (const DistanceConstraint &distanceConstraint : lengthConstraints){
setNeighbors(&distanceConstraint, {
distanceConstraint.a,
distanceConstraint.b
});
}
#pragma omp parallel for default(none) shared(setNeighbors, volumeConstraints)
for (const VolumeConstraint &volumeConstraint : volumeConstraints){
setNeighbors(&volumeConstraint, {
volumeConstraint.a,
volumeConstraint.b,
volumeConstraint.c,
volumeConstraint.d,
});
}
vector<unordered_set<const Constraint *>> partitions = createPartitions(graph);
constraintData.partitionCount = partitions.size();
reorderConstraintIndices(partitions);
}
void SoftBody::reorderConstraintIndices(const vector<unordered_set<const Constraint *>> &partitions) {
ConstraintData reordered;
for (uint32_t partition = 0; partition < partitions.size(); partition++){
reordered.recordNewPartition();
for (const Constraint * constraint : partitions[partition])
constraint->writeData(reordered);
reordered.writePartitionInformation();
}
std::swap(reordered.edgePartitions, constraintData.edgePartitions);
std::swap(reordered.tetrahedronPartitions, constraintData.tetrahedronPartitions);
std::swap(reordered.edges, constraintData.edges);
std::swap(reordered.tetrahedra, constraintData.tetrahedra);
//std::swap(reordered, constraintData);
//reordered.triangles.swap(constraintData.triangles);
//reordered.tetrahedra.swap(constraintData.tetrahedra);
//std::swap(reordered.partitionCount, constraintData.partitionCount);
} }

2
src/vulkan/descriptor_pool.cpp
Unescape View File

@ -49,7 +49,7 @@ DescriptorPool::DescriptorPool() {
// sizes // sizes
addBinding(DescriptorSet::MESH, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, VK_SHADER_STAGE_COMPUTE_BIT); addBinding(DescriptorSet::MESH, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, VK_SHADER_STAGE_COMPUTE_BIT);
// dt // simulation properties
addBinding(DescriptorSet::SIMULATION, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, VK_SHADER_STAGE_COMPUTE_BIT); addBinding(DescriptorSet::SIMULATION, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER, VK_SHADER_STAGE_COMPUTE_BIT);
for (const auto &[set, bindings] : setBindings) for (const auto &[set, bindings] : setBindings)

Write Preview
Loading…
Cancel
Save