Adding SDF LevelSet #187

elalish · 2022-08-18T20:16:16Z

Fixes #129

I've implemented my own take on Marching Cubes for making manifold meshes from signed-distance functions. It's a form of Marching Tetrahedra on a body-centered cubic grid, which ensures all tetrahedra are identical, maximum quality, and symmetric. It also ensures all output is manifold without needing to deal with edge cases.

I've implemented this for the GPU (of course it works on the CPU as well) because the computation can get pretty intense. The interface is intended to mimic Thrust with its use of functors.

These new tests also unearthed an existing and very tricky bug, whose fix is also included here.

This reverts commit ba2348f.

fire · 2022-08-18T20:34:09Z

So if I have a random glb I may be possible to test this?

Do I need to generate sdfs in some form? What form is that?

elalish · 2022-08-19T05:26:17Z

So if I have a random glb I may be possible to test this?

Do I need to generate sdfs in some form? What form is that?

There's documentation in sdf.h for the SDF class. Any function that takes a vec3 and returns a float will do. Turning an arbitrary GLB into an SDF is pretty non-trivial - that's not really what it's for. It's about creating complex surfaces from mathematical functions. Try out the sample.

pca006132

Some comments about readability and use of atomics. Haven't yet look into details of the algorithmic part.

pca006132 · 2022-08-19T15:01:40Z

src/sdf/include/sdf.h

+#else
+  std::atomic<Uint64>& tar = reinterpret_cast<std::atomic<Uint64>&>(target);
+  Uint64 old_val = tar.load();
+  tar.compare_exchange_weak(compare, val);


compare_exchange may fail. It will fail when there are interleaved CAS (the latter one will fail), and will fail spuriously when compare_exchange_weak is used (like when the cache line is touched, as it can be implemented with better efficiency in some architecture like arm with exclusive monitor).

Generally, when you just want to update the value atomically, you have to do it in a loop, repeat the CAS until the return value is true (updated successfully).

Thanks, I'm still a little fuzzy on atomics. I've been impressed how fast they are in CUDA even when hit pretty often. I'll change that for the CPU side.

Yeah it seems that the implementation of atomic in CUDA is different from that of CPU. Actually, the slow thing about atomic is about cache update across cores, I guess this is a bit different in GPU.

Thank you, I've just read about this now. It sounds likes weak means I can't rely on return == compare to guarantee a swap has taken place. I could probably tweak it to make use of the boolean return instead, but I think there's no big problem going with strong.

pca006132 · 2022-08-19T15:03:12Z

src/sdf/include/sdf.h

+}
+
+__host__ __device__ Uint64 SpreadBits3(Uint64 v) {
+  v = v & 0x1fffff;


These magic numbers look really magical, it would be nice if you can explain how/where did you get them.

I have similar ones in the collider; Morton codes are pretty standard and various versions of the magic numbers can be found in algorithm reference books.

pca006132 · 2022-08-19T15:13:39Z

src/sdf/include/sdf.h

+    uint32_t idx = vert.key & (Size() - 1);
+    while (1) {
+      Uint64& key = table_[idx].key;
+      const Uint64 found = AtomicCAS(key, kOpen, vert.key);


Regarding the implementation of a concurrent hashtable, I'm a bit worried about the performance of having atomic operations in a loop. Atomic operations are a lot slower than typical instructions, and the cost is still a lot higher even when you only have 1 thread running.

If writes are not common, I think a lock might give you better performance. If writes are common, I think I can have a look and see if there is a way to reduce the use of atomic operations.

As usual I was focused more on the GPU, where locks aren't much of an option. I'm totally open to recommendations though.

You can implement a spin lock with atomic.

My impression was that atomics are mostly slow when concurrent access is actually happening; do they have much performance impact otherwise? I would expect very little concurrency on these - they will only happen during a hash collision (which is already rare) that also happens to be concurrent. I made the step large so the nearby threads don't tend to step on each other.

Yes, they have performance impact even when there is only 1 thread running these atomic instructions.

That's good to know, thanks. I'll leave it this way for now, but you're welcome to optimize it - that will be a good learning experience for me.

src/sdf/include/sdf.h

pca006132 · 2022-08-19T15:21:33Z

src/sdf/include/sdf.h

+  inline __host__ __device__ void operator()(Uint64 mortonCode) {
+    const glm::ivec4 gridIndex = DecodeMorton(mortonCode);
+
+    if (glm::any(glm::greaterThan(glm::ivec3(gridIndex), gridSize))) return;


Also, just noticed that you have some glm::any(glm::some_comparison(...)) patterns. Maybe we can define some helper functions to make it easier to read? Like anyGreaterThan or something.

That's potentially a lot of permutations, I think I'm going to leave that for now.

src/sdf/include/sdf.h

pca006132 · 2022-08-20T15:27:27Z

src/sdf/include/sdf.h

+    gridVert.distance = BoundedSDF(gridIndex);
+
+    bool keep = false;
+    float minDist2 = 0.25 * 0.25;


What is the usage of minDist2?

Ah yes, that's dead code now. I had a cool idea, but realized it would create edge cases that could break manifoldness.

pca006132 · 2022-08-20T15:50:22Z

src/sdf/include/sdf.h

+    tet[2] = thisVert.Inside();
+    int edges[6] = {base.edgeVerts[0], -1, -1, -1, -1, -1};
+    for (const int i : {0, 1, 2}) {
+      edges[1] = base.edgeVerts[i + 1];


This part is relatively complicated, maybe we can reorder some statements and avoid reusing the array to make it more readable?

for (const int i : {0, 1, 2}) { thisIndex = leadIndex; thisIndex[Prev3(i)] -= 1; GridVert nextVert = gridVerts[MortonCode(thisIndex)]; tet[3] = nextVert.Inside(); int edges1[6] = { base.edgeVerts[0], base.edgeVerts[i+1], nextVert.edgeVerts[Next3(i) + 4], nextVert.edgeVerts[Prev3(i) + 1], thisVert.edgeVerts[i + 4], base.edgeVerts[Prev3(i) + 4] }; thisVert = nextVert; if (!skipTet && thisVert.key != kOpen) CreateTris(tet, edges1); skipTet = thisVert.key == kOpen; tet[2] = tet[3]; int edges2[6] = ... }

(I personally think that it is easy to get lost reading many array element operations, especially with some order dependencies with other variables)

Okay, this has been cleaned up a lot; thanks for the suggestion, I ended up using that and a few other things too.

codecov-commenter · 2022-08-27T17:50:34Z

Codecov Report

Merging #187 (94a3c69) into master (08598c0) will increase coverage by 0.02%.
The diff coverage is 97.35%.

@@            Coverage Diff             @@
##           master     #187      +/-   ##
==========================================
+ Coverage   98.15%   98.18%   +0.02%     
==========================================
  Files          30       33       +3     
  Lines        2602     2919     +317     
==========================================
+ Hits         2554     2866     +312     
- Misses         48       53       +5

Impacted Files	Coverage Δ
src/collider/src/collider.cpp	`100.00% <ø> (ø)`
src/manifold/src/impl.cpp	`100.00% <ø> (ø)`
src/utilities/include/structs.h	`100.00% <ø> (ø)`
src/utilities/include/utils.h	`100.00% <ø> (ø)`
src/sdf/include/sdf.h	`96.10% <96.10%> (ø)`
samples/src/gyroid_module.cpp	`100.00% <100.00%> (ø)`
src/manifold/src/boolean3.cpp	`100.00% <100.00%> (ø)`
src/manifold/src/boolean_result.cpp	`99.57% <100.00%> (+0.06%)`	⬆️
src/sdf/src/sdf.cpp	`100.00% <100.00%> (ø)`
src/utilities/include/par.h	`100.00% <100.00%> (ø)`
... and 17 more

Help us with your feedback. Take ten seconds to tell us how you rate us. Have a feature suggestion? Share it here.

elalish · 2022-08-30T20:05:09Z

I've simplified the API; I think this is pretty much ready to merge.

* polymorphic device functions working * starting to flesh out algorithm * getting closer * fleshed out vert computation * more progress * BuildTris * fixed some syntax * updating to cpp * added VecDc * SDF compiles * SDF runs * fixed some bugs * manifold output * fixed out of memory bug * flattened sides * back to egg-crate * works with CUDA too * moved SDF * Revert "moved SDF" This reverts commit ba2348f. * added gyroid sample * added sdf_test * added SDF tests * tests pass * updated gyroid_module * fixed gyroid module * added docs * addressing feedback * fix cuda detection * clean up detect cuda * fixed CUDA failure * consistent execution policy in Boolean * remove SDF wrapper class * simplified tet logic * use NeighborInside() * optimized SDF * moved RemoveUnreferencedVerts * put printf statements behind MANIFOLD_DEBUG * enable hashtable resize * added test for hashtable resizing

elalish added 30 commits March 11, 2022 11:45

polymorphic device functions working

11716ba

starting to flesh out algorithm

e34b939

getting closer

c869c50

fleshed out vert computation

677f7b8

merging master

4ef8f42

more progress

055035f

BuildTris

5efb432

fixed some syntax

0dd1f57

merging master

c17e9cb

updating to cpp

0859717

Merge branch 'master' of github.com:elalish/manifold into sdf

06d1a38

added VecDc

0280a40

Merge branch 'master' of github.com:elalish/manifold into sdf

6936c56

SDF compiles

09cb0cf

SDF runs

3888973

fixed some bugs

245338a

manifold output

997fb73

fixed out of memory bug

f4eeca7

flattened sides

941fd62

merging master

f0f997a

merging master

25f376c

Merge branch 'master' of github.com:elalish/manifold into sdf

d5e4635

back to egg-crate

6dd429a

works with CUDA too

2c95950

moved SDF

ba2348f

Revert "moved SDF"

059f86d

This reverts commit ba2348f.

added gyroid sample

8b1a17e

added sdf_test

73e395d

added SDF tests

0fc8c8f

tests pass

01bdb70

elalish added 3 commits August 18, 2022 11:59

updated gyroid_module

fbcb707

fixed gyroid module

ea0c7d0

added docs

7a57b0f

elalish self-assigned this Aug 18, 2022

elalish requested a review from pca006132 August 18, 2022 20:16

pca006132 reviewed Aug 19, 2022

View reviewed changes

pca006132 reviewed Aug 20, 2022

View reviewed changes

elalish added 4 commits August 22, 2022 23:09

addressing feedback

cd57ffe

fix cuda detection

c1f825f

clean up detect cuda

4d89088

fixed CUDA failure

a2b282f

elalish mentioned this pull request Aug 27, 2022

dynamic backend implementation #126

Merged

2 tasks

elalish added 2 commits August 27, 2022 10:13

consistent execution policy in Boolean

b91db71

remove SDF wrapper class

94a3c69

elalish added 7 commits August 27, 2022 11:24

simplified tet logic

c626f4e

use NeighborInside()

55de478

optimized SDF

1ba085e

moved RemoveUnreferencedVerts

0aefceb

put printf statements behind MANIFOLD_DEBUG

3ccbba9

enable hashtable resize

c56f1bc

added test for hashtable resizing

7bdaa27

elalish requested a review from pca006132 August 30, 2022 20:03

elalish changed the title ~~Adding SDF class~~ Adding SDF LevelSet Sep 1, 2022

elalish merged commit 11e58ea into master Sep 1, 2022

elalish deleted the sdf branch September 1, 2022 16:04

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Adding SDF LevelSet #187

Adding SDF LevelSet #187

elalish commented Aug 18, 2022 •

edited

Loading

fire commented Aug 18, 2022 •

edited

Loading

elalish commented Aug 19, 2022

pca006132 left a comment

pca006132 Aug 19, 2022

elalish Aug 20, 2022

pca006132 Aug 20, 2022

elalish Aug 23, 2022

pca006132 Aug 19, 2022

elalish Aug 30, 2022

pca006132 Aug 19, 2022

elalish Aug 20, 2022

pca006132 Aug 20, 2022

elalish Aug 23, 2022

pca006132 Aug 23, 2022

elalish Aug 27, 2022

pca006132 Aug 19, 2022

elalish Aug 30, 2022

pca006132 Aug 20, 2022

elalish Aug 23, 2022

pca006132 Aug 20, 2022

pca006132 Aug 20, 2022

elalish Aug 30, 2022

codecov-commenter commented Aug 27, 2022

elalish commented Aug 30, 2022

Adding SDF LevelSet #187

Adding SDF LevelSet #187

Conversation

elalish commented Aug 18, 2022 • edited Loading

fire commented Aug 18, 2022 • edited Loading

elalish commented Aug 19, 2022

pca006132 left a comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

Choose a reason for hiding this comment

codecov-commenter commented Aug 27, 2022

Codecov Report

elalish commented Aug 30, 2022

elalish commented Aug 18, 2022 •

edited

Loading

fire commented Aug 18, 2022 •

edited

Loading