Skip to content

Commit

Permalink
added task shader support
Browse files Browse the repository at this point in the history
  • Loading branch information
@Try
Try committed Oct 30, 2024
1 parent 99e6f3a commit 30b0b98
Show file tree
Hide file tree
Showing 16 changed files with 304 additions and 87 deletions.
33 changes: 20 additions & 13 deletions reference/opt/shaders-msl/mesh/mesh-shader-basic-lines.msl3.spv14.vk.nocompat.mesh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -89,6 +89,13 @@ struct BlockOutPrim
float4 b;
};

struct TaskPayload
{
float a;
float b;
int c;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 3u, 4u);

struct spvPerVertex
Expand Down Expand Up @@ -116,33 +123,33 @@ struct spvPerPrimitive
using spvMesh_t = mesh<spvPerVertex, spvPerPrimitive, 24, 22, topology::line>;

static inline __attribute__((always_inline))
void _4(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, thread uint& gl_LocalInvocationIndex, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, threadgroup uint2& spvMeshSizes)
void _4(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, thread uint& gl_LocalInvocationIndex, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, const object_data TaskPayload& payload, threadgroup uint2& spvMeshSizes)
{
spvSetMeshOutputsEXT(gl_LocalInvocationIndex, spvMeshSizes, 24u, 22u);
float3 _158 = float3(gl_GlobalInvocationID);
float _159 = _158.x;
gl_MeshVerticesEXT[gl_LocalInvocationIndex].gl_Position = float4(_159, _158.yz, 1.0);
float3 _163 = float3(gl_GlobalInvocationID);
float _164 = _163.x;
gl_MeshVerticesEXT[gl_LocalInvocationIndex].gl_Position = float4(_164, _163.yz, 1.0);
gl_MeshVerticesEXT[gl_LocalInvocationIndex].gl_PointSize = 2.0;
gl_MeshVerticesEXT[gl_LocalInvocationIndex].gl_ClipDistance[0] = 4.0;
vOut[gl_LocalInvocationIndex] = float4(_159, _158.yz, 2.0);
vOut[gl_LocalInvocationIndex] = float4(_164, _163.yz, 2.0);
outputs[gl_LocalInvocationIndex].a = float4(5.0);
outputs[gl_LocalInvocationIndex].b = float4(6.0);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (gl_LocalInvocationIndex < 22u)
{
vPrim[gl_LocalInvocationIndex] = float4(float3(gl_WorkGroupID), 3.0);
prim_outputs[gl_LocalInvocationIndex].a = float4(0.0);
prim_outputs[gl_LocalInvocationIndex].b = float4(1.0);
prim_outputs[gl_LocalInvocationIndex].a = float4(payload.a);
prim_outputs[gl_LocalInvocationIndex].b = float4(payload.b);
gl_PrimitiveLineIndicesEXT[gl_LocalInvocationIndex] = uint2(0u, 1u) + uint2(gl_LocalInvocationIndex);
int _206 = int(gl_GlobalInvocationID.x);
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_PrimitiveID = _206;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_Layer = _206 + 1;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_ViewportIndex = _206 + 2;
int _217 = int(gl_GlobalInvocationID.x);
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_PrimitiveID = _217;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_Layer = _217 + 1;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_ViewportIndex = _217 + 2;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_CullPrimitiveEXT = short((gl_GlobalInvocationID.x & 1u) != 0u);
}
}

[[mesh]] void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], spvMesh_t spvMesh)
[[mesh]] void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], spvMesh_t spvMesh, const object_data TaskPayload& payload [[payload]])
{
threadgroup uint2 spvMeshSizes;
threadgroup spvUnsafeArray<uint2, 22> gl_PrimitiveLineIndicesEXT;
Expand All @@ -154,7 +161,7 @@ void _4(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, threa
threadgroup spvUnsafeArray<BlockOutPrim, 22> prim_outputs;
threadgroup spvUnsafeArray<float, 16> shared_float;
if (gl_LocalInvocationIndex == 0) spvMeshSizes.y = 0u;
_4(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID, gl_MeshVerticesEXT, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, spvMeshSizes);
_4(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID, gl_MeshVerticesEXT, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, payload, spvMeshSizes);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (spvMeshSizes.y == 0)
{
Expand Down
32 changes: 16 additions & 16 deletions reference/opt/shaders-msl/mesh/mesh-shader-basic-triangle.msl3.spv14.vk.nocompat.mesh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -81,6 +81,13 @@ struct BlockOutPrim
float4 b;
};

struct TaskPayload
{
float a;
float b;
int c;
};

struct gl_MeshPerPrimitiveEXT
{
uint gl_PrimitiveID [[primitive_id]];
Expand All @@ -91,13 +98,6 @@ struct gl_MeshPerPrimitiveEXT

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 3u, 4u);

struct TaskPayload
{
float a;
float b;
int c;
};

struct spvPerVertex
{
float4 gl_Position [[position]];
Expand All @@ -123,7 +123,7 @@ struct spvPerPrimitive
using spvMesh_t = mesh<spvPerVertex, spvPerPrimitive, 24, 22, topology::triangle>;

static inline __attribute__((always_inline))
void _4(threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, thread uint& gl_LocalInvocationIndex, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, threadgroup spvUnsafeArray<uint3, 22>& gl_PrimitiveTriangleIndicesEXT, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, threadgroup uint2& spvMeshSizes)
void _4(threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, thread uint& gl_LocalInvocationIndex, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, const object_data TaskPayload& payload, threadgroup spvUnsafeArray<uint3, 22>& gl_PrimitiveTriangleIndicesEXT, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, threadgroup uint2& spvMeshSizes)
{
spvSetMeshOutputsEXT(gl_LocalInvocationIndex, spvMeshSizes, 24u, 22u);
float3 _27 = float3(gl_GlobalInvocationID);
Expand All @@ -138,18 +138,18 @@ void _4(threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT,
if (gl_LocalInvocationIndex < 22u)
{
vPrim[gl_LocalInvocationIndex] = float4(float3(gl_WorkGroupID), 3.0);
prim_outputs[gl_LocalInvocationIndex].a = float4(0.0);
prim_outputs[gl_LocalInvocationIndex].b = float4(1.0);
prim_outputs[gl_LocalInvocationIndex].a = float4(payload.a);
prim_outputs[gl_LocalInvocationIndex].b = float4(payload.b);
gl_PrimitiveTriangleIndicesEXT[gl_LocalInvocationIndex] = uint3(0u, 1u, 2u) + uint3(gl_LocalInvocationIndex);
int _116 = int(gl_GlobalInvocationID.x);
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_PrimitiveID = _116;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_Layer = _116 + 1;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_ViewportIndex = _116 + 2;
int _123 = int(gl_GlobalInvocationID.x);
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_PrimitiveID = _123;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_Layer = _123 + 1;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_ViewportIndex = _123 + 2;
gl_MeshPrimitivesEXT[gl_LocalInvocationIndex].gl_CullPrimitiveEXT = short((gl_GlobalInvocationID.x & 1u) != 0u);
}
}

[[mesh]] void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], spvMesh_t spvMesh)
[[mesh]] void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], spvMesh_t spvMesh, const object_data TaskPayload& payload [[payload]])
{
threadgroup uint2 spvMeshSizes;
threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24> gl_MeshVerticesEXT;
Expand All @@ -161,7 +161,7 @@ void _4(threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT,
threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22> gl_MeshPrimitivesEXT;
threadgroup spvUnsafeArray<float, 16> shared_float;
if (gl_LocalInvocationIndex == 0) spvMeshSizes.y = 0u;
_4(gl_MeshVerticesEXT, gl_LocalInvocationIndex, gl_GlobalInvocationID, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, gl_PrimitiveTriangleIndicesEXT, gl_MeshPrimitivesEXT, spvMeshSizes);
_4(gl_MeshVerticesEXT, gl_LocalInvocationIndex, gl_GlobalInvocationID, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, payload, gl_PrimitiveTriangleIndicesEXT, gl_MeshPrimitivesEXT, spvMeshSizes);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (spvMeshSizes.y == 0)
{
Expand Down
22 changes: 22 additions & 0 deletions reference/opt/shaders-msl/task/task-basic.msl3.spv14.vk.nocompat.task
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,22 @@
#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct TaskPayload
{
float a;
float b;
int c;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);

[[object]] void main0(mesh_grid_properties spvMgp, object_data TaskPayload& payload [[payload]])
{
payload.a = 1.2000000476837158203125;
payload.b = 2.2999999523162841796875;
payload.c = 3;
spvMgp.set_threadgroups_per_grid(uint3(1u, 2u, 3u));
}

23 changes: 23 additions & 0 deletions reference/opt/shaders-msl/task/task-func.msl3.spv14.vk.nocompat.task
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,23 @@
#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

struct TaskPayload
{
float a;
float b;
int c;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);

[[object]] void main0(mesh_grid_properties spvMgp, object_data TaskPayload& payload [[payload]])
{
payload.a = 1.2000000476837158203125;
payload.b = 2.2999999523162841796875;
payload.c = 3;
threadgroup_barrier(mem_flags::mem_threadgroup);
spvMgp.set_threadgroups_per_grid(uint3(1u, 2u, 3u));
}

18 changes: 9 additions & 9 deletions reference/shaders-msl/mesh/mesh-shader-basic-lines.msl3.spv14.vk.nocompat.mesh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -89,15 +89,15 @@ struct BlockOutPrim
float4 b;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 3u, 4u);

struct TaskPayload
{
float a;
float b;
int c;
};

constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(2u, 3u, 4u);

struct spvPerVertex
{
float4 gl_Position [[position]];
Expand Down Expand Up @@ -133,7 +133,7 @@ void main3(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, th
}

static inline __attribute__((always_inline))
void main2(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, thread uint& gl_LocalInvocationIndex, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, threadgroup uint2& spvMeshSizes)
void main2(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, thread uint& gl_LocalInvocationIndex, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, const object_data TaskPayload& payload, threadgroup uint2& spvMeshSizes)
{
spvSetMeshOutputsEXT(gl_LocalInvocationIndex, spvMeshSizes, 24u, 22u);
gl_MeshVerticesEXT[gl_LocalInvocationIndex].gl_Position = float4(float3(gl_GlobalInvocationID), 1.0);
Expand All @@ -146,19 +146,19 @@ void main2(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, th
if (gl_LocalInvocationIndex < 22u)
{
vPrim[gl_LocalInvocationIndex] = float4(float3(gl_WorkGroupID), 3.0);
prim_outputs[gl_LocalInvocationIndex].a = float4(0.0);
prim_outputs[gl_LocalInvocationIndex].b = float4(1.0);
prim_outputs[gl_LocalInvocationIndex].a = float4(payload.a);
prim_outputs[gl_LocalInvocationIndex].b = float4(payload.b);
main3(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID);
}
}

static inline __attribute__((always_inline))
void _4(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, thread uint& gl_LocalInvocationIndex, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, threadgroup uint2& spvMeshSizes)
void _4(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, thread uint& gl_LocalInvocationIndex, threadgroup spvUnsafeArray<gl_MeshPerPrimitiveEXT, 22>& gl_MeshPrimitivesEXT, thread uint3& gl_GlobalInvocationID, threadgroup spvUnsafeArray<gl_MeshPerVertexEXT, 24>& gl_MeshVerticesEXT, threadgroup spvUnsafeArray<float4, 24>& vOut, threadgroup spvUnsafeArray<BlockOut, 24>& outputs, threadgroup spvUnsafeArray<float4, 22>& vPrim, thread uint3& gl_WorkGroupID, threadgroup spvUnsafeArray<BlockOutPrim, 22>& prim_outputs, const object_data TaskPayload& payload, threadgroup uint2& spvMeshSizes)
{
main2(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID, gl_MeshVerticesEXT, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, spvMeshSizes);
main2(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID, gl_MeshVerticesEXT, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, payload, spvMeshSizes);
}

[[mesh]] void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], spvMesh_t spvMesh)
[[mesh]] void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], spvMesh_t spvMesh, const object_data TaskPayload& payload [[payload]])
{
threadgroup uint2 spvMeshSizes;
threadgroup spvUnsafeArray<uint2, 22> gl_PrimitiveLineIndicesEXT;
Expand All @@ -170,7 +170,7 @@ void _4(threadgroup spvUnsafeArray<uint2, 22>& gl_PrimitiveLineIndicesEXT, threa
threadgroup spvUnsafeArray<BlockOutPrim, 22> prim_outputs;
threadgroup spvUnsafeArray<float, 16> shared_float;
if (gl_LocalInvocationIndex == 0) spvMeshSizes.y = 0u;
_4(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID, gl_MeshVerticesEXT, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, spvMeshSizes);
_4(gl_PrimitiveLineIndicesEXT, gl_LocalInvocationIndex, gl_MeshPrimitivesEXT, gl_GlobalInvocationID, gl_MeshVerticesEXT, vOut, outputs, vPrim, gl_WorkGroupID, prim_outputs, payload, spvMeshSizes);
threadgroup_barrier(mem_flags::mem_threadgroup);
if (spvMeshSizes.y == 0)
{
Expand Down
Loading

0 comments on commit 30b0b98

Please sign in to comment.