mirror of
https://github.com/italicsjenga/vello.git
synced 2025-01-10 12:41:30 +11:00
Add generated shaders
This commit is contained in:
parent
4ed339d434
commit
b36ca7fc2e
72
tests/shader/gen/prefix_reduce.hlsl
Normal file
72
tests/shader/gen/prefix_reduce.hlsl
Normal file
|
@ -0,0 +1,72 @@
|
||||||
|
struct Monoid
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
|
||||||
|
|
||||||
|
ByteAddressBuffer _40 : register(t0);
|
||||||
|
RWByteAddressBuffer _129 : register(u1);
|
||||||
|
|
||||||
|
static uint3 gl_WorkGroupID;
|
||||||
|
static uint3 gl_LocalInvocationID;
|
||||||
|
static uint3 gl_GlobalInvocationID;
|
||||||
|
struct SPIRV_Cross_Input
|
||||||
|
{
|
||||||
|
uint3 gl_WorkGroupID : SV_GroupID;
|
||||||
|
uint3 gl_LocalInvocationID : SV_GroupThreadID;
|
||||||
|
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
|
||||||
|
};
|
||||||
|
|
||||||
|
groupshared Monoid sh_scratch[512];
|
||||||
|
|
||||||
|
Monoid combine_monoid(Monoid a, Monoid b)
|
||||||
|
{
|
||||||
|
Monoid _22 = { a.element + b.element };
|
||||||
|
return _22;
|
||||||
|
}
|
||||||
|
|
||||||
|
void comp_main()
|
||||||
|
{
|
||||||
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
||||||
|
Monoid _44;
|
||||||
|
_44.element = _40.Load(ix * 4 + 0);
|
||||||
|
Monoid agg;
|
||||||
|
agg.element = _44.element;
|
||||||
|
Monoid param_1;
|
||||||
|
for (uint i = 1u; i < 8u; i++)
|
||||||
|
{
|
||||||
|
Monoid param = agg;
|
||||||
|
Monoid _64;
|
||||||
|
_64.element = _40.Load((ix + i) * 4 + 0);
|
||||||
|
param_1.element = _64.element;
|
||||||
|
agg = combine_monoid(param, param_1);
|
||||||
|
}
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
for (uint i_1 = 0u; i_1 < 9u; i_1++)
|
||||||
|
{
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u)
|
||||||
|
{
|
||||||
|
Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))];
|
||||||
|
Monoid param_2 = agg;
|
||||||
|
Monoid param_3 = other;
|
||||||
|
agg = combine_monoid(param_2, param_3);
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
}
|
||||||
|
if (gl_LocalInvocationID.x == 0u)
|
||||||
|
{
|
||||||
|
_129.Store(gl_WorkGroupID.x * 4 + 0, agg.element);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
[numthreads(512, 1, 1)]
|
||||||
|
void main(SPIRV_Cross_Input stage_input)
|
||||||
|
{
|
||||||
|
gl_WorkGroupID = stage_input.gl_WorkGroupID;
|
||||||
|
gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
|
||||||
|
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
|
||||||
|
comp_main();
|
||||||
|
}
|
68
tests/shader/gen/prefix_reduce.msl
Normal file
68
tests/shader/gen/prefix_reduce.msl
Normal file
|
@ -0,0 +1,68 @@
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
struct Monoid
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Monoid_1
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct InBuf
|
||||||
|
{
|
||||||
|
Monoid_1 inbuf[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
struct OutBuf
|
||||||
|
{
|
||||||
|
Monoid_1 outbuf[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
|
||||||
|
|
||||||
|
static inline __attribute__((always_inline))
|
||||||
|
Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b)
|
||||||
|
{
|
||||||
|
return Monoid{ a.element + b.element };
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void main0(const device InBuf& _40 [[buffer(0)]], device OutBuf& _129 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
||||||
|
{
|
||||||
|
threadgroup Monoid sh_scratch[512];
|
||||||
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
||||||
|
Monoid agg;
|
||||||
|
agg.element = _40.inbuf[ix].element;
|
||||||
|
Monoid param_1;
|
||||||
|
for (uint i = 1u; i < 8u; i++)
|
||||||
|
{
|
||||||
|
Monoid param = agg;
|
||||||
|
param_1.element = _40.inbuf[ix + i].element;
|
||||||
|
agg = combine_monoid(param, param_1);
|
||||||
|
}
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
for (uint i_1 = 0u; i_1 < 9u; i_1++)
|
||||||
|
{
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
if ((gl_LocalInvocationID.x + uint(1 << int(i_1))) < 512u)
|
||||||
|
{
|
||||||
|
Monoid other = sh_scratch[gl_LocalInvocationID.x + uint(1 << int(i_1))];
|
||||||
|
Monoid param_2 = agg;
|
||||||
|
Monoid param_3 = other;
|
||||||
|
agg = combine_monoid(param_2, param_3);
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
}
|
||||||
|
if (gl_LocalInvocationID.x == 0u)
|
||||||
|
{
|
||||||
|
_129.outbuf[gl_WorkGroupID.x].element = agg.element;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
BIN
tests/shader/gen/prefix_reduce.spv
Normal file
BIN
tests/shader/gen/prefix_reduce.spv
Normal file
Binary file not shown.
80
tests/shader/gen/prefix_root.hlsl
Normal file
80
tests/shader/gen/prefix_root.hlsl
Normal file
|
@ -0,0 +1,80 @@
|
||||||
|
struct Monoid
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
|
||||||
|
|
||||||
|
static const Monoid _133 = { 0u };
|
||||||
|
|
||||||
|
RWByteAddressBuffer _42 : register(u0);
|
||||||
|
|
||||||
|
static uint3 gl_LocalInvocationID;
|
||||||
|
static uint3 gl_GlobalInvocationID;
|
||||||
|
struct SPIRV_Cross_Input
|
||||||
|
{
|
||||||
|
uint3 gl_LocalInvocationID : SV_GroupThreadID;
|
||||||
|
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
|
||||||
|
};
|
||||||
|
|
||||||
|
groupshared Monoid sh_scratch[512];
|
||||||
|
|
||||||
|
Monoid combine_monoid(Monoid a, Monoid b)
|
||||||
|
{
|
||||||
|
Monoid _22 = { a.element + b.element };
|
||||||
|
return _22;
|
||||||
|
}
|
||||||
|
|
||||||
|
void comp_main()
|
||||||
|
{
|
||||||
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
||||||
|
Monoid _46;
|
||||||
|
_46.element = _42.Load(ix * 4 + 0);
|
||||||
|
Monoid local[8];
|
||||||
|
local[0].element = _46.element;
|
||||||
|
Monoid param_1;
|
||||||
|
for (uint i = 1u; i < 8u; i++)
|
||||||
|
{
|
||||||
|
Monoid param = local[i - 1u];
|
||||||
|
Monoid _71;
|
||||||
|
_71.element = _42.Load((ix + i) * 4 + 0);
|
||||||
|
param_1.element = _71.element;
|
||||||
|
local[i] = combine_monoid(param, param_1);
|
||||||
|
}
|
||||||
|
Monoid agg = local[7];
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
for (uint i_1 = 0u; i_1 < 9u; i_1++)
|
||||||
|
{
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
|
||||||
|
{
|
||||||
|
Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
|
||||||
|
Monoid param_2 = other;
|
||||||
|
Monoid param_3 = agg;
|
||||||
|
agg = combine_monoid(param_2, param_3);
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
Monoid row = _133;
|
||||||
|
if (gl_LocalInvocationID.x > 0u)
|
||||||
|
{
|
||||||
|
row = sh_scratch[gl_LocalInvocationID.x - 1u];
|
||||||
|
}
|
||||||
|
for (uint i_2 = 0u; i_2 < 8u; i_2++)
|
||||||
|
{
|
||||||
|
Monoid param_4 = row;
|
||||||
|
Monoid param_5 = local[i_2];
|
||||||
|
Monoid m = combine_monoid(param_4, param_5);
|
||||||
|
_42.Store((ix + i_2) * 4 + 0, m.element);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
[numthreads(512, 1, 1)]
|
||||||
|
void main(SPIRV_Cross_Input stage_input)
|
||||||
|
{
|
||||||
|
gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
|
||||||
|
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
|
||||||
|
comp_main();
|
||||||
|
}
|
112
tests/shader/gen/prefix_root.msl
Normal file
112
tests/shader/gen/prefix_root.msl
Normal file
|
@ -0,0 +1,112 @@
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-braces"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
template<typename T, size_t Num>
|
||||||
|
struct spvUnsafeArray
|
||||||
|
{
|
||||||
|
T elements[Num ? Num : 1];
|
||||||
|
|
||||||
|
thread T& operator [] (size_t pos) thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const thread T& operator [] (size_t pos) const thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
device T& operator [] (size_t pos) device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const device T& operator [] (size_t pos) const device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
constexpr const constant T& operator [] (size_t pos) const constant
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
threadgroup T& operator [] (size_t pos) threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Monoid
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Monoid_1
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct DataBuf
|
||||||
|
{
|
||||||
|
Monoid_1 data[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
|
||||||
|
|
||||||
|
static inline __attribute__((always_inline))
|
||||||
|
Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b)
|
||||||
|
{
|
||||||
|
return Monoid{ a.element + b.element };
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void main0(device DataBuf& _42 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
|
||||||
|
{
|
||||||
|
threadgroup Monoid sh_scratch[512];
|
||||||
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
||||||
|
spvUnsafeArray<Monoid, 8> local;
|
||||||
|
local[0].element = _42.data[ix].element;
|
||||||
|
Monoid param_1;
|
||||||
|
for (uint i = 1u; i < 8u; i++)
|
||||||
|
{
|
||||||
|
Monoid param = local[i - 1u];
|
||||||
|
param_1.element = _42.data[ix + i].element;
|
||||||
|
local[i] = combine_monoid(param, param_1);
|
||||||
|
}
|
||||||
|
Monoid agg = local[7];
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
for (uint i_1 = 0u; i_1 < 9u; i_1++)
|
||||||
|
{
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
|
||||||
|
{
|
||||||
|
Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
|
||||||
|
Monoid param_2 = other;
|
||||||
|
Monoid param_3 = agg;
|
||||||
|
agg = combine_monoid(param_2, param_3);
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
Monoid row = Monoid{ 0u };
|
||||||
|
if (gl_LocalInvocationID.x > 0u)
|
||||||
|
{
|
||||||
|
row = sh_scratch[gl_LocalInvocationID.x - 1u];
|
||||||
|
}
|
||||||
|
for (uint i_2 = 0u; i_2 < 8u; i_2++)
|
||||||
|
{
|
||||||
|
Monoid param_4 = row;
|
||||||
|
Monoid param_5 = local[i_2];
|
||||||
|
Monoid m = combine_monoid(param_4, param_5);
|
||||||
|
_42.data[ix + i_2].element = m.element;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
BIN
tests/shader/gen/prefix_root.spv
Normal file
BIN
tests/shader/gen/prefix_root.spv
Normal file
Binary file not shown.
92
tests/shader/gen/prefix_scan.hlsl
Normal file
92
tests/shader/gen/prefix_scan.hlsl
Normal file
|
@ -0,0 +1,92 @@
|
||||||
|
struct Monoid
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
static const uint3 gl_WorkGroupSize = uint3(512u, 1u, 1u);
|
||||||
|
|
||||||
|
static const Monoid _133 = { 0u };
|
||||||
|
|
||||||
|
RWByteAddressBuffer _42 : register(u0);
|
||||||
|
RWByteAddressBuffer _143 : register(u1);
|
||||||
|
|
||||||
|
static uint3 gl_WorkGroupID;
|
||||||
|
static uint3 gl_LocalInvocationID;
|
||||||
|
static uint3 gl_GlobalInvocationID;
|
||||||
|
struct SPIRV_Cross_Input
|
||||||
|
{
|
||||||
|
uint3 gl_WorkGroupID : SV_GroupID;
|
||||||
|
uint3 gl_LocalInvocationID : SV_GroupThreadID;
|
||||||
|
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
|
||||||
|
};
|
||||||
|
|
||||||
|
groupshared Monoid sh_scratch[512];
|
||||||
|
|
||||||
|
Monoid combine_monoid(Monoid a, Monoid b)
|
||||||
|
{
|
||||||
|
Monoid _22 = { a.element + b.element };
|
||||||
|
return _22;
|
||||||
|
}
|
||||||
|
|
||||||
|
void comp_main()
|
||||||
|
{
|
||||||
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
||||||
|
Monoid _46;
|
||||||
|
_46.element = _42.Load(ix * 4 + 0);
|
||||||
|
Monoid local[8];
|
||||||
|
local[0].element = _46.element;
|
||||||
|
Monoid param_1;
|
||||||
|
for (uint i = 1u; i < 8u; i++)
|
||||||
|
{
|
||||||
|
Monoid param = local[i - 1u];
|
||||||
|
Monoid _71;
|
||||||
|
_71.element = _42.Load((ix + i) * 4 + 0);
|
||||||
|
param_1.element = _71.element;
|
||||||
|
local[i] = combine_monoid(param, param_1);
|
||||||
|
}
|
||||||
|
Monoid agg = local[7];
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
for (uint i_1 = 0u; i_1 < 9u; i_1++)
|
||||||
|
{
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
|
||||||
|
{
|
||||||
|
Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
|
||||||
|
Monoid param_2 = other;
|
||||||
|
Monoid param_3 = agg;
|
||||||
|
agg = combine_monoid(param_2, param_3);
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
}
|
||||||
|
GroupMemoryBarrierWithGroupSync();
|
||||||
|
Monoid row = _133;
|
||||||
|
if (gl_WorkGroupID.x > 0u)
|
||||||
|
{
|
||||||
|
Monoid _148;
|
||||||
|
_148.element = _143.Load((gl_WorkGroupID.x - 1u) * 4 + 0);
|
||||||
|
row.element = _148.element;
|
||||||
|
}
|
||||||
|
if (gl_LocalInvocationID.x > 0u)
|
||||||
|
{
|
||||||
|
Monoid param_4 = row;
|
||||||
|
Monoid param_5 = sh_scratch[gl_LocalInvocationID.x - 1u];
|
||||||
|
row = combine_monoid(param_4, param_5);
|
||||||
|
}
|
||||||
|
for (uint i_2 = 0u; i_2 < 8u; i_2++)
|
||||||
|
{
|
||||||
|
Monoid param_6 = row;
|
||||||
|
Monoid param_7 = local[i_2];
|
||||||
|
Monoid m = combine_monoid(param_6, param_7);
|
||||||
|
_42.Store((ix + i_2) * 4 + 0, m.element);
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
||||||
|
[numthreads(512, 1, 1)]
|
||||||
|
void main(SPIRV_Cross_Input stage_input)
|
||||||
|
{
|
||||||
|
gl_WorkGroupID = stage_input.gl_WorkGroupID;
|
||||||
|
gl_LocalInvocationID = stage_input.gl_LocalInvocationID;
|
||||||
|
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
|
||||||
|
comp_main();
|
||||||
|
}
|
123
tests/shader/gen/prefix_scan.msl
Normal file
123
tests/shader/gen/prefix_scan.msl
Normal file
|
@ -0,0 +1,123 @@
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-braces"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
template<typename T, size_t Num>
|
||||||
|
struct spvUnsafeArray
|
||||||
|
{
|
||||||
|
T elements[Num ? Num : 1];
|
||||||
|
|
||||||
|
thread T& operator [] (size_t pos) thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const thread T& operator [] (size_t pos) const thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
device T& operator [] (size_t pos) device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const device T& operator [] (size_t pos) const device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
constexpr const constant T& operator [] (size_t pos) const constant
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
threadgroup T& operator [] (size_t pos) threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Monoid
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct Monoid_1
|
||||||
|
{
|
||||||
|
uint element;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct DataBuf
|
||||||
|
{
|
||||||
|
Monoid_1 data[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
struct ParentBuf
|
||||||
|
{
|
||||||
|
Monoid_1 parent[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(512u, 1u, 1u);
|
||||||
|
|
||||||
|
static inline __attribute__((always_inline))
|
||||||
|
Monoid combine_monoid(thread const Monoid& a, thread const Monoid& b)
|
||||||
|
{
|
||||||
|
return Monoid{ a.element + b.element };
|
||||||
|
}
|
||||||
|
|
||||||
|
kernel void main0(device DataBuf& _42 [[buffer(0)]], device ParentBuf& _143 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
||||||
|
{
|
||||||
|
threadgroup Monoid sh_scratch[512];
|
||||||
|
uint ix = gl_GlobalInvocationID.x * 8u;
|
||||||
|
spvUnsafeArray<Monoid, 8> local;
|
||||||
|
local[0].element = _42.data[ix].element;
|
||||||
|
Monoid param_1;
|
||||||
|
for (uint i = 1u; i < 8u; i++)
|
||||||
|
{
|
||||||
|
Monoid param = local[i - 1u];
|
||||||
|
param_1.element = _42.data[ix + i].element;
|
||||||
|
local[i] = combine_monoid(param, param_1);
|
||||||
|
}
|
||||||
|
Monoid agg = local[7];
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
for (uint i_1 = 0u; i_1 < 9u; i_1++)
|
||||||
|
{
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
if (gl_LocalInvocationID.x >= uint(1 << int(i_1)))
|
||||||
|
{
|
||||||
|
Monoid other = sh_scratch[gl_LocalInvocationID.x - uint(1 << int(i_1))];
|
||||||
|
Monoid param_2 = other;
|
||||||
|
Monoid param_3 = agg;
|
||||||
|
agg = combine_monoid(param_2, param_3);
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
sh_scratch[gl_LocalInvocationID.x] = agg;
|
||||||
|
}
|
||||||
|
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||||
|
Monoid row = Monoid{ 0u };
|
||||||
|
if (gl_WorkGroupID.x > 0u)
|
||||||
|
{
|
||||||
|
row.element = _143.parent[gl_WorkGroupID.x - 1u].element;
|
||||||
|
}
|
||||||
|
if (gl_LocalInvocationID.x > 0u)
|
||||||
|
{
|
||||||
|
Monoid param_4 = row;
|
||||||
|
Monoid param_5 = sh_scratch[gl_LocalInvocationID.x - 1u];
|
||||||
|
row = combine_monoid(param_4, param_5);
|
||||||
|
}
|
||||||
|
for (uint i_2 = 0u; i_2 < 8u; i_2++)
|
||||||
|
{
|
||||||
|
Monoid param_6 = row;
|
||||||
|
Monoid param_7 = local[i_2];
|
||||||
|
Monoid m = combine_monoid(param_6, param_7);
|
||||||
|
_42.data[ix + i_2].element = m.element;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
BIN
tests/shader/gen/prefix_scan.spv
Normal file
BIN
tests/shader/gen/prefix_scan.spv
Normal file
Binary file not shown.
Loading…
Reference in a new issue