Support XFB in MoltenVK #2169

gpx1000 · 2023-06-19T23:44:35Z

Adds support for MoltenVK to use TransformFeedback

This is analogous to the existing support for fixing up shader inputs. It is intended to be used with tessellation to add implicit builtins that are read from a later stage, despite not being written in an earlier stage. (Believe it or not, this is in fact legal in Vulkan.) Helps fix 8 CTS tests under `dEQP-VK.pipeline.*.no_position`. (Eight other tests work solely by accident without this change.)

MSL: Add a mechanism to fix up shader outputs. Approved-by: Steven Winston

It is possible in SPIR-V to declare multiple specialization constants with the same constant ID. The most common cause of this in GLSL is defining a spec constant, then declaring the workgroup size to use that spec constant by its ID. But, MSL forbids defining multiple function constants with the same function constant ID. So, we must only emit one definition of the actual function constant (with the `[[function_constant(id)]]` attribute); but we can point the other variables at this one definition. Fixes three tests in the Vulkan CTS under `dEQP-VK.compute.basic.max_local_size_*`.

MSL: Deduplicate function constants. Approved-by: Steven Winston

Does analysis of outputs and sorts them into buffers. Nothing else yet.

# Conflicts: # main.cpp # spirv_cross_c.cpp # spirv_cross_c.h # spirv_msl.cpp # spirv_msl.hpp

This only does the bare minimum needed to write XFB data (and not even that actually). It still needs to calculate the offset in the buffer where the data need to be written, and primitive types other than points need to be implemented.

# Conflicts: # reference/shaders-msl/comp/local-size-duplicate-spec-id.comp # spirv_msl.cpp

…e updated to be used.

HansKristian-Work · 2023-07-03T11:46:17Z

Given this is still marked draft, I assume this is still WIP and I'll be notified when it's ready for review?

…ernels.

…sform feedback.

This lets us reference it later.

I don't expect this to build, let alone work. (Really, all these changes ought to be squashed when merged to SPIRV-Cross.)

Work out how indexing works for triangle fans. A little bit closer...

This is the only type of shader that can even have such outputs. Not only does this save some work in most cases, it also fixes a problem with the next patch.

We still rely on it to pass around and collect the output. To avoid duplicates, only do this if we would not do this normally.

Use the offset from the `XfbOutput` struct instead of querying it again from the ID. Use the `member_index` local instead of using `size() - 1` when setting member decorations. Don't set the qualified name for builtin block variabless--we handle those a different way. Use the member index from the `XfbOutput` when inspecing the original block type instead of the `member_index` local. This one was a real bug; honestly, I don't know how it even worked before.

…types. Make sure they use the `thread` AS and that they have the `packed_` prefix, if necessary.

Add missing changes from previous patch.

...on the builder and not for me.

…hash<enum> is supposed to just work.

…ordered_map.

…et also doesn't support enum in C++11 for hash key

HansKristian-Work · 2023-09-25T08:43:20Z

reference/opt/shaders-msl/vert/transform-feedback-decorations.xfb-line-list.vert

+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _20 = {};
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))
+        return;


This breaks threadgroup_barrier.

HansKristian-Work · 2023-09-25T08:44:34Z

reference/opt/shaders-msl/vert/transform-feedback-decorations.xfb-point-list.vert

+    if (all(gl_GlobalInvocationID.xy == 0))
+    {
+        uint spvWritten = spvStageInputSize.x * spvStageInputSize.y;
+        atomic_store_explicit(spvXfbCounter1, spvInitOffset1 + sizeof(*spvXfb1) * spvWritten, memory_order_relaxed);


How is XFB ordering maintained here? XFB data must be emitted in-order with input primitives.

The actual XFB buffers are indexed by the global invocation ID.

HansKristian-Work · 2023-09-25T08:56:32Z

reference/shaders-msl/vert/transform-feedback-decorations.xfb-triangle-list.vert

+    spvXfb3 = reinterpret_cast<device spvXfbBuffer3*>(reinterpret_cast<device char*>(spvXfb3) + spvInitOffset3);
+    if ((gl_GlobalInvocationID.x % 3u == 2) || gl_GlobalInvocationID.x + 2 < spvStageInputSize.x)
+        spvXfb3[spvXfbIndex] = spvXfbOutput3;
+    threadgroup_barrier(mem_flags::mem_device);


Why is a barrier even needed here?

To make sure the atomic load of the counter happens-before the atomic update, particularly since Metal doesn't support acquire/release semantics for atomics.

A different workgroup can touch the counter here though, so not sure threadgroup_barrier is enough to ensure the in-order requirement of XFB.

Only one thread in the entire dispatch is allowed to write the counter.

...But if the first workgroup completes before the others, they may wind up loading the counter after the first thread writes to it. Hmm, this may be tougher than I thought...

The only reasonable solution is to manage the counter from the outside and only pass the offset (read-only) into the shader.

HansKristian-Work · 2023-09-25T08:57:53Z

reference/shaders-msl/vert/transform-feedback-decorations.xfb-triangle-strip.vert

+    spvXfbBuffer1 spvXfbOutput1 = {};
+    spvXfbBuffer2 spvXfbOutput2 = {};
+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _25 = {};


Where does vertex output go?

Directly to the transform feedback buffers.

How is this rasterized? Do you re-shade through the normal vertex shader in a second pass, or do you shade from the XFB feedback as well?

We intend to have MoltenVK shade from the XFB data.

I think that is broken. Some reasons why:

It's possible to only capture a subset of varyings.

Robustness rules. XFB buffers can be exhausted, where you still rasterize, but you must stop writing XFB data.

HansKristian-Work · 2023-09-25T08:58:30Z

reference/shaders-msl/vert/transform-feedback-decorations.xfb-triangle-strip.vert

+    spvXfbBuffer2 spvXfbOutput2 = {};
+    spvXfbBuffer3 spvXfbOutput3 = {};
+    VertOut _25 = {};
+    if (any(gl_GlobalInvocationID >= spvStageInputSize))


What about index buffers?

HansKristian-Work

First of all, this is 58 (!) commits. I'm not reviewing that as-is. Please rebase this into something more digestable.

Also, given the scope of this, please add a PR description that explains the implementation strategy, problem scenarios, which corner cases don't work, etc, also why this feature is even desirable in the first place. Does it pass CTS?

Only gets the base of the primitive so far.

cdavis5e and others added 15 commits August 23, 2022 15:26

Merged in msl-shader-output-fixup (pull request KhronosGroup#2)

2127a3b

MSL: Add a mechanism to fix up shader outputs. Approved-by: Steven Winston

Merged in msl-duplicate-spec-id (pull request KhronosGroup#8)

f195855

MSL: Deduplicate function constants. Approved-by: Steven Winston

Checkpoint for transform feedback work.

343ff6e

Does analysis of outputs and sorts them into buffers. Nothing else yet.

Merge remote-tracking branch 'origin/master'

048ac2d

# Conflicts: # main.cpp # spirv_cross_c.cpp # spirv_cross_c.h # spirv_msl.cpp # spirv_msl.hpp

Get things building.

179c6e0

Checkpoint: Beginnings of writing XFB data.

f1c0ad2

This only does the bare minimum needed to write XFB data (and not even that actually). It still needs to calculate the offset in the buffer where the data need to be written, and primitive types other than points need to be implemented.

get it building.

117eaa3

Merge branch 'master' into xfb

1e8cbe4

# Conflicts: # reference/shaders-msl/comp/local-size-duplicate-spec-id.comp # spirv_msl.cpp

get xfb decorations shader to work.

f1913aa

check in for direction adjustment.

f8a27d9

Dynamic is an undefined primitive type. xfb_primitive_type needs to b…

cebb964

…e updated to be used.

Working together with Chip

37c0972

fix warnings from CI

aab161a

gpx1000 marked this pull request as draft June 20, 2023 00:17

Merge branch 'master' into origin-xfb

9d2329a

cdavis5e added 12 commits July 11, 2023 16:07

Merge remote-tracking branch 'origin/main' into xfb

111cebb

Merge remote-tracking branch 'origin/main' into xfb

e3cf900

Make sure vertex functions that use transform feedback become Metal k…

562b959

…ernels.

Merge remote-tracking branch 'gpx1000/xfb' into xfb

d62fe77

Add the transform feedback buffer parameters to the vertex shader.

36d39df

Make sure all used outputs, including builtins, get XFB buffers.

35858fb

Make sure builtins have the correct names in XFB buffers.

3c427de

Add command line parameter to set the primitive type assumed for tran…

c352f94

…sform feedback.

Add a variable for the XFB counter buffer.

001ff7d

This lets us reference it later.

Really crappy checkpoint for XFB work.

28babde

I don't expect this to build, let alone work. (Really, all these changes ought to be squashed when merged to SPIRV-Cross.)

Getting closer...

fb520f4

Fix indices of triangle strips to account for winding.

556c9fa

Work out how indexing works for triangle fans. A little bit closer...

cdavis5e added 10 commits September 17, 2023 23:41

Fix broken constant generated MSL.

16dd1f1

Only create a per-patch output block for tessellation control shaders.

742f725

This is the only type of shader that can even have such outputs. Not only does this save some work in most cases, it also fixes a problem with the next patch.

Make sure the local variable for an output block gets created.

8f66f30

We still rely on it to pass around and collect the output. To avoid duplicates, only do this if we would not do this normally.

Make sure captured outputs passed as implicit arguments have correct …

8dbf250

…types. Make sure they use the `thread` AS and that they have the `packed_` prefix, if necessary.

Only use qualified name for builtins in the entry point().

b020270

Add missing changes from previous patch.

Remove extraneous right parentheses.

3521814

Clang-format the changes.

bf4f823

Add tests for transform feedback in MSL.

109959e

Merge branch 'main' into xfb

3bd855f

cdavis5e marked this pull request as ready for review September 19, 2023 01:18

cdavis5e requested a review from HansKristian-Work September 19, 2023 01:18

cdavis5e and others added 8 commits September 18, 2023 22:03

Merge remote-tracking branch 'origin' into xfb

1154932

Merge remote-tracking branch 'steve/xfb' into xfb

a547b52

Attempt to fix MSVC build.

64fa0b6

Try again to work around MSVC brokenness.

dada588

Attempt to work around weird brokenness that only happens...

adb3a7b

...on the builder and not for me.

Try again to get the stupid compiler on the builder to see that std::…

fec7607

…hash<enum> is supposed to just work.

Testing hypothesis that C++11 doesn't support enum as a key for an un…

0393302

…ordered_map.

hypothesis was correct for unordered_map stands to reason unordered_s…

739a140

…et also doesn't support enum in C++11 for hash key

HansKristian-Work reviewed Sep 25, 2023

View reviewed changes

HansKristian-Work requested changes Sep 25, 2023

View reviewed changes

Merge remote-tracking branch 'origin' into xfb

15a8b70

cdavis5e marked this pull request as draft November 23, 2023 03:49

cdavis5e added 2 commits November 28, 2023 22:17

Merge remote-tracking branch 'origin' into xfb

575e75d

Unfinished support for XFB+tessellation.

8bcfd32

Only gets the base of the primitive so far.

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Support XFB in MoltenVK #2169

Support XFB in MoltenVK #2169

gpx1000 commented Jun 19, 2023

HansKristian-Work commented Jul 3, 2023

HansKristian-Work Sep 25, 2023

HansKristian-Work Sep 25, 2023

cdavis5e Sep 25, 2023

HansKristian-Work Sep 25, 2023

cdavis5e Sep 25, 2023

HansKristian-Work Sep 26, 2023

cdavis5e Sep 26, 2023

HansKristian-Work Sep 27, 2023

HansKristian-Work Sep 25, 2023

cdavis5e Sep 25, 2023

HansKristian-Work Sep 26, 2023

cdavis5e Sep 26, 2023

HansKristian-Work Oct 5, 2023

HansKristian-Work Sep 25, 2023

HansKristian-Work left a comment •

edited

Loading

Support XFB in MoltenVK #2169

Are you sure you want to change the base?

Support XFB in MoltenVK #2169

Conversation

gpx1000 commented Jun 19, 2023

HansKristian-Work commented Jul 3, 2023

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

HansKristian-Work left a comment • edited Loading

Choose a reason for hiding this comment

HansKristian-Work left a comment •

edited

Loading