Skip to content

Commit 5af8a04

Browse files
HansKristian-WorkHansKristian-Work
authored
Merge pull request #1148 from KhronosGroup/dynamic-offsets
Merge dynamic offset support with some bug-fixes.
2 parents 537bee3 + afa5480 commit 5af8a04

12 files changed

Lines changed: 385 additions & 8 deletions

‎CMakeLists.txt‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -287,7 +287,7 @@ if (SPIRV_CROSS_STATIC)
287287
endif()
288288

289289
set(spirv-cross-abi-major 0)
290-
set(spirv-cross-abi-minor 17)
290+
set(spirv-cross-abi-minor 18)
291291
set(spirv-cross-abi-patch 0)
292292

293293
if (SPIRV_CROSS_SHARED)

‎main.cpp‎

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -522,6 +522,7 @@ struct CLIArguments
522522
bool vulkan_glsl_disable_ext_samplerless_texture_functions = false;
523523
bool emit_line_directives = false;
524524
SmallVector<uint32_t> msl_discrete_descriptor_sets;
525+
SmallVector<pair<uint32_t, uint32_t>> msl_dynamic_buffers;
525526
SmallVector<PLSArg> pls_in;
526527
SmallVector<PLSArg> pls_out;
527528
SmallVector<Remap> remaps;
@@ -600,6 +601,7 @@ static void print_help()
600601
"\t[--msl-multiview]\n"
601602
"\t[--msl-view-index-from-device-index]\n"
602603
"\t[--msl-dispatch-base]\n"
604+
"\t[--msl-dynamic-buffer <set index> <binding>]\n"
603605
"\t[--hlsl]\n"
604606
"\t[--reflect]\n"
605607
"\t[--shader-model]\n"
@@ -764,6 +766,9 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
764766
msl_comp->set_msl_options(msl_opts);
765767
for (auto &v : args.msl_discrete_descriptor_sets)
766768
msl_comp->add_discrete_descriptor_set(v);
769+
uint32_t i = 0;
770+
for (auto &v : args.msl_dynamic_buffers)
771+
msl_comp->add_dynamic_buffer(v.first, v.second, i++);
767772
}
768773
else if (args.hlsl)
769774
compiler.reset(new CompilerHLSL(move(spirv_parser.get_parsed_ir())));
@@ -1086,6 +1091,13 @@ static int main_inner(int argc, char *argv[])
10861091
cbs.add("--msl-view-index-from-device-index",
10871092
[&args](CLIParser &) { args.msl_view_index_from_device_index = true; });
10881093
cbs.add("--msl-dispatch-base", [&args](CLIParser &) { args.msl_dispatch_base = true; });
1094+
cbs.add("--msl-dynamic-buffer", [&args](CLIParser &parser) {
1095+
args.msl_argument_buffers = true;
1096+
// Make sure next_uint() is called in-order.
1097+
uint32_t desc_set = parser.next_uint();
1098+
uint32_t binding = parser.next_uint();
1099+
args.msl_dynamic_buffers.push_back(make_pair(desc_set, binding));
1100+
});
10891101
cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
10901102
cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
10911103
auto old_name = parser.next_string();

‎reference/opt/shaders-msl/comp/basic.dynamic-buffer.msl2.comp‎

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
#include <metal_stdlib>
2+
#include <simd/simd.h>
3+
4+
using namespace metal;
5+
6+
struct Baz
7+
{
8+
int e;
9+
int f;
10+
};
11+
12+
struct Foo
13+
{
14+
int a;
15+
int b;
16+
};
17+
18+
struct Bar
19+
{
20+
int c;
21+
int d;
22+
};
23+
24+
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u);
25+
26+
struct spvDescriptorSetBuffer0
27+
{
28+
constant Foo* m_34 [[id(0)]];
29+
constant Bar* m_40 [[id(1)]];
30+
};
31+
32+
struct spvDescriptorSetBuffer1
33+
{
34+
device Baz* baz [[id(0)]][3][3][2];
35+
};
36+
37+
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]])
38+
{
39+
constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]);
40+
device Baz* baz[3][3][2] =
41+
{
42+
{
43+
{
44+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][0] + spvDynamicOffsets[1]),
45+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][1] + spvDynamicOffsets[2]),
46+
},
47+
{
48+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][0] + spvDynamicOffsets[3]),
49+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][1] + spvDynamicOffsets[4]),
50+
},
51+
{
52+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][0] + spvDynamicOffsets[5]),
53+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][1] + spvDynamicOffsets[6]),
54+
},
55+
},
56+
{
57+
{
58+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][0] + spvDynamicOffsets[7]),
59+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][1] + spvDynamicOffsets[8]),
60+
},
61+
{
62+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][0] + spvDynamicOffsets[9]),
63+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][1] + spvDynamicOffsets[10]),
64+
},
65+
{
66+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][0] + spvDynamicOffsets[11]),
67+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][1] + spvDynamicOffsets[12]),
68+
},
69+
},
70+
{
71+
{
72+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][0] + spvDynamicOffsets[13]),
73+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][1] + spvDynamicOffsets[14]),
74+
},
75+
{
76+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][0] + spvDynamicOffsets[15]),
77+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][1] + spvDynamicOffsets[16]),
78+
},
79+
{
80+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][0] + spvDynamicOffsets[17]),
81+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][1] + spvDynamicOffsets[18]),
82+
},
83+
},
84+
};
85+
86+
baz[gl_GlobalInvocationID.x][gl_GlobalInvocationID.y][gl_GlobalInvocationID.z]->e = _34.a + (*spvDescriptorSet0.m_40).c;
87+
baz[gl_GlobalInvocationID.x][gl_GlobalInvocationID.y][gl_GlobalInvocationID.z]->f = _34.b * (*spvDescriptorSet0.m_40).d;
88+
}
89+

‎reference/shaders-msl/comp/basic.dynamic-buffer.msl2.comp‎

Lines changed: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
1+
#include <metal_stdlib>
2+
#include <simd/simd.h>
3+
4+
using namespace metal;
5+
6+
struct Baz
7+
{
8+
int e;
9+
int f;
10+
};
11+
12+
struct Foo
13+
{
14+
int a;
15+
int b;
16+
};
17+
18+
struct Bar
19+
{
20+
int c;
21+
int d;
22+
};
23+
24+
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(3u, 3u, 2u);
25+
26+
struct spvDescriptorSetBuffer0
27+
{
28+
constant Foo* m_34 [[id(0)]];
29+
constant Bar* m_40 [[id(1)]];
30+
};
31+
32+
struct spvDescriptorSetBuffer1
33+
{
34+
device Baz* baz [[id(0)]][3][3][2];
35+
};
36+
37+
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]])
38+
{
39+
constant auto& _34 = *(constant Foo* )((constant char* )spvDescriptorSet0.m_34 + spvDynamicOffsets[0]);
40+
device Baz* baz[3][3][2] =
41+
{
42+
{
43+
{
44+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][0] + spvDynamicOffsets[1]),
45+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][0][1] + spvDynamicOffsets[2]),
46+
},
47+
{
48+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][0] + spvDynamicOffsets[3]),
49+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][1][1] + spvDynamicOffsets[4]),
50+
},
51+
{
52+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][0] + spvDynamicOffsets[5]),
53+
(device Baz* )((device char* )spvDescriptorSet1.baz[0][2][1] + spvDynamicOffsets[6]),
54+
},
55+
},
56+
{
57+
{
58+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][0] + spvDynamicOffsets[7]),
59+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][0][1] + spvDynamicOffsets[8]),
60+
},
61+
{
62+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][0] + spvDynamicOffsets[9]),
63+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][1][1] + spvDynamicOffsets[10]),
64+
},
65+
{
66+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][0] + spvDynamicOffsets[11]),
67+
(device Baz* )((device char* )spvDescriptorSet1.baz[1][2][1] + spvDynamicOffsets[12]),
68+
},
69+
},
70+
{
71+
{
72+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][0] + spvDynamicOffsets[13]),
73+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][0][1] + spvDynamicOffsets[14]),
74+
},
75+
{
76+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][0] + spvDynamicOffsets[15]),
77+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][1][1] + spvDynamicOffsets[16]),
78+
},
79+
{
80+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][0] + spvDynamicOffsets[17]),
81+
(device Baz* )((device char* )spvDescriptorSet1.baz[2][2][1] + spvDynamicOffsets[18]),
82+
},
83+
},
84+
};
85+
86+
uint3 coords = gl_GlobalInvocationID;
87+
baz[coords.x][coords.y][coords.z]->e = _34.a + (*spvDescriptorSet0.m_40).c;
88+
baz[coords.x][coords.y][coords.z]->f = _34.b * (*spvDescriptorSet0.m_40).d;
89+
}
90+

‎shaders-msl/comp/basic.dynamic-buffer.msl2.comp‎

Lines changed: 27 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,27 @@
1+
#version 450
2+
layout(local_size_x = 3, local_size_y = 3, local_size_z = 2) in;
3+
4+
layout(set = 0, binding = 0) uniform Foo
5+
{
6+
int a;
7+
int b;
8+
};
9+
10+
layout(set = 0, binding = 1) uniform Bar
11+
{
12+
int c;
13+
int d;
14+
};
15+
16+
layout(set = 1, binding = 2) buffer Baz
17+
{
18+
int e;
19+
int f;
20+
} baz[3][3][2];
21+
22+
void main()
23+
{
24+
uvec3 coords = gl_GlobalInvocationID;
25+
baz[coords.x][coords.y][coords.z].e = a + c;
26+
baz[coords.x][coords.y][coords.z].f = b * d;
27+
}

‎spirv_cross_c.cpp‎

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -565,6 +565,10 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
565565
case SPVC_COMPILER_OPTION_MSL_DISPATCH_BASE:
566566
options->msl.dispatch_base = value != 0;
567567
break;
568+
569+
case SPVC_COMPILER_OPTION_MSL_DYNAMIC_OFFSETS_BUFFER_INDEX:
570+
options->msl.dynamic_offsets_buffer_index = value;
571+
break;
568572
#endif
569573

570574
default:
@@ -902,6 +906,27 @@ spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler,
902906
#endif
903907
}
904908

909+
spvc_result spvc_compiler_msl_add_dynamic_buffer(spvc_compiler compiler, unsigned desc_set, unsigned binding, unsigned index)
910+
{
911+
#if SPIRV_CROSS_C_API_MSL
912+
if (compiler->backend != SPVC_BACKEND_MSL)
913+
{
914+
compiler->context->report_error("MSL function used on a non-MSL backend.");
915+
return SPVC_ERROR_INVALID_ARGUMENT;
916+
}
917+
918+
auto &msl = *static_cast<CompilerMSL *>(compiler->compiler.get());
919+
msl.add_dynamic_buffer(desc_set, binding, index);
920+
return SPVC_SUCCESS;
921+
#else
922+
(void)binding;
923+
(void)desc_set;
924+
(void)index;
925+
compiler->context->report_error("MSL function used on a non-MSL backend.");
926+
return SPVC_ERROR_INVALID_ARGUMENT;
927+
#endif
928+
}
929+
905930
spvc_result spvc_compiler_msl_add_discrete_descriptor_set(spvc_compiler compiler, unsigned desc_set)
906931
{
907932
#if SPIRV_CROSS_C_API_MSL

‎spirv_cross_c.h‎

Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -33,7 +33,7 @@ extern "C" {
3333
/* Bumped if ABI or API breaks backwards compatibility. */
3434
#define SPVC_C_API_VERSION_MAJOR 0
3535
/* Bumped if APIs or enumerations are added in a backwards compatible way. */
36-
#define SPVC_C_API_VERSION_MINOR 17
36+
#define SPVC_C_API_VERSION_MINOR 18
3737
/* Bumped if internal implementation details change. */
3838
#define SPVC_C_API_VERSION_PATCH 0
3939

@@ -526,6 +526,7 @@ typedef enum spvc_compiler_option
526526
SPVC_COMPILER_OPTION_MSL_DEVICE_INDEX = 40 | SPVC_COMPILER_OPTION_MSL_BIT,
527527
SPVC_COMPILER_OPTION_MSL_VIEW_INDEX_FROM_DEVICE_INDEX = 41 | SPVC_COMPILER_OPTION_MSL_BIT,
528528
SPVC_COMPILER_OPTION_MSL_DISPATCH_BASE = 42 | SPVC_COMPILER_OPTION_MSL_BIT,
529+
SPVC_COMPILER_OPTION_MSL_DYNAMIC_OFFSETS_BUFFER_INDEX = 43 | SPVC_COMPILER_OPTION_MSL_BIT,
529530

530531
SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
531532
} spvc_compiler_option;
@@ -632,6 +633,8 @@ SPVC_PUBLIC_API spvc_result spvc_compiler_msl_set_fragment_output_components(spv
632633
SPVC_PUBLIC_API unsigned spvc_compiler_msl_get_automatic_resource_binding(spvc_compiler compiler, spvc_variable_id id);
633634
SPVC_PUBLIC_API unsigned spvc_compiler_msl_get_automatic_resource_binding_secondary(spvc_compiler compiler, spvc_variable_id id);
634635

636+
SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_dynamic_buffer(spvc_compiler compiler, unsigned desc_set, unsigned binding, unsigned index);
637+
635638
/*
636639
* Reflect resources.
637640
* Maps almost 1:1 to C++ API.

‎spirv_glsl.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -12483,6 +12483,14 @@ void CompilerGLSL::end_scope()
1248312483
statement("}");
1248412484
}
1248512485

12486+
void CompilerGLSL::end_scope(const string &trailer)
12487+
{
12488+
if (!indent)
12489+
SPIRV_CROSS_THROW("Popping empty indent stack.");
12490+
indent--;
12491+
statement("}", trailer);
12492+
}
12493+
1248612494
void CompilerGLSL::end_scope_decl()
1248712495
{
1248812496
if (!indent)

‎spirv_glsl.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -332,6 +332,7 @@ class CompilerGLSL : public Compiler
332332

333333
void begin_scope();
334334
void end_scope();
335+
void end_scope(const std::string &trailer);
335336
void end_scope_decl();
336337
void end_scope_decl(const std::string &decl);
337338

0 commit comments

Comments
 (0)