cb35934248
Vulkan has two types of buffer descriptors, `VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER_DYNAMIC` and `VK_DESCRIPTOR_TYPE_STORAGE_BUFFER_DYNAMIC`, which allow the client to offset the buffers by an amount given when the descriptor set is bound to a pipeline. Metal provides no direct support for this when the buffer in question is in an argument buffer, so once again we're on our own. These offsets cannot be stored or associated in any way with the argument buffer itself, because they are set at bind time. Different pipelines may have different offsets set. Therefore, we must use a separate buffer, not in any argument buffer, to hold these offsets. Then the shader must manually offset the buffer pointer. This change fully supports arrays, including arrays of arrays, even though Vulkan forbids them. It does not, however, support runtime arrays. Perhaps later.
90 lines
3.2 KiB
Plaintext
90 lines
3.2 KiB
Plaintext
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using namespace metal;
|
|
|
|
struct Baz
|
|
{
|
|
int e;
|
|
int f;
|
|
};
|
|
|
|
struct Foo
|
|
{
|
|
int a;
|
|
int b;
|
|
};
|
|
|
|
struct Bar
|
|
{
|
|
int c;
|
|
int d;
|
|
};
|
|
|
|
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u);
|
|
|
|
struct spvDescriptorSetBuffer0
|
|
{
|
|
constant Foo* m_34 [[id(0)]];
|
|
constant Bar* m_40 [[id(1)]];
|
|
};
|
|
|
|
struct spvDescriptorSetBuffer1
|
|
{
|
|
device Baz* baz [[id(0)]][3][3][2];
|
|
};
|
|
|
|
kernel void main0(constant spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], constant spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], constant uint* spvDynamicOffsets [[buffer(23)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
|
{
|
|
constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]);
|
|
device Baz* baz[3][3][2] =
|
|
{
|
|
{
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][0] + spvDynamicOffsets[1]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][1] + spvDynamicOffsets[2]),
|
|
},
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][0] + spvDynamicOffsets[3]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][1] + spvDynamicOffsets[4]),
|
|
},
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][0] + spvDynamicOffsets[5]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][1] + spvDynamicOffsets[6]),
|
|
},
|
|
},
|
|
{
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][0] + spvDynamicOffsets[7]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][1] + spvDynamicOffsets[8]),
|
|
},
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][0] + spvDynamicOffsets[9]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][1] + spvDynamicOffsets[10]),
|
|
},
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][0] + spvDynamicOffsets[11]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][1] + spvDynamicOffsets[12]),
|
|
},
|
|
},
|
|
{
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][0] + spvDynamicOffsets[13]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][1] + spvDynamicOffsets[14]),
|
|
},
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][0] + spvDynamicOffsets[15]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][1] + spvDynamicOffsets[16]),
|
|
},
|
|
{
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][0] + spvDynamicOffsets[17]),
|
|
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][1] + spvDynamicOffsets[18]),
|
|
},
|
|
},
|
|
};
|
|
|
|
baz[gl_GlobalInvocationID.x][gl_GlobalInvocationID.y][gl_GlobalInvocationID.z]->e = _34.a + (*spvDescriptorSet0.m_40).c;
|
|
baz[gl_GlobalInvocationID.x][gl_GlobalInvocationID.y][gl_GlobalInvocationID.z]->f = _34.b * (*spvDescriptorSet0.m_40).d;
|
|
}
|
|
|