Skip to content

Commit

Permalink
Updated spirv-cross.
Browse files Browse the repository at this point in the history
  • Loading branch information
bkaradzic committed Sep 7, 2019
1 parent f3a0654 commit 3dd517b
Show file tree
Hide file tree
Showing 87 changed files with 5,658 additions and 1,205 deletions.
16 changes: 15 additions & 1 deletion 3rdparty/spirv-cross/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -287,7 +287,7 @@ if (SPIRV_CROSS_STATIC)
endif()

set(spirv-cross-abi-major 0)
set(spirv-cross-abi-minor 16)
set(spirv-cross-abi-minor 18)
set(spirv-cross-abi-patch 0)

if (SPIRV_CROSS_SHARED)
Expand Down Expand Up @@ -461,6 +461,14 @@ if (SPIRV_CROSS_CLI)
target_link_libraries(spirv-cross-msl-resource-binding-test spirv-cross-c)
set_target_properties(spirv-cross-msl-resource-binding-test PROPERTIES LINK_FLAGS "${spirv-cross-link-flags}")

add_executable(spirv-cross-msl-ycbcr-conversion-test tests-other/msl_ycbcr_conversion_test.cpp)
target_link_libraries(spirv-cross-msl-ycbcr-conversion-test spirv-cross-c)
set_target_properties(spirv-cross-msl-ycbcr-conversion-test PROPERTIES LINK_FLAGS "${spirv-cross-link-flags}")

add_executable(spirv-cross-typed-id-test tests-other/typed_id_test.cpp)
target_link_libraries(spirv-cross-typed-id-test spirv-cross-core)
set_target_properties(spirv-cross-typed-id-test PROPERTIES LINK_FLAGS "${spirv-cross-link-flags}")

if (CMAKE_COMPILER_IS_GNUCXX OR (${CMAKE_CXX_COMPILER_ID} MATCHES "Clang"))
target_compile_options(spirv-cross-c-api-test PRIVATE -std=c89 -Wall -Wextra)
endif()
Expand All @@ -475,6 +483,12 @@ if (SPIRV_CROSS_CLI)
COMMAND $<TARGET_FILE:spirv-cross-msl-constexpr-test> ${CMAKE_CURRENT_SOURCE_DIR}/tests-other/msl_constexpr_test.spv)
add_test(NAME spirv-cross-msl-resource-binding-test
COMMAND $<TARGET_FILE:spirv-cross-msl-resource-binding-test> ${CMAKE_CURRENT_SOURCE_DIR}/tests-other/msl_resource_binding.spv)
add_test(NAME spirv-cross-msl-ycbcr-conversion-test
COMMAND $<TARGET_FILE:spirv-cross-msl-ycbcr-conversion-test> ${CMAKE_CURRENT_SOURCE_DIR}/tests-other/msl_ycbcr_conversion_test.spv)
add_test(NAME spirv-cross-msl-ycbcr-conversion-test-2
COMMAND $<TARGET_FILE:spirv-cross-msl-ycbcr-conversion-test> ${CMAKE_CURRENT_SOURCE_DIR}/tests-other/msl_ycbcr_conversion_test_2.spv)
add_test(NAME spirv-cross-typed-id-test
COMMAND $<TARGET_FILE:spirv-cross-typed-id-test>)
add_test(NAME spirv-cross-test
COMMAND ${PYTHON_EXECUTABLE} ${CMAKE_CURRENT_SOURCE_DIR}/test_shaders.py --parallel
${spirv-cross-externals}
Expand Down
18 changes: 15 additions & 3 deletions 3rdparty/spirv-cross/main.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -246,7 +246,7 @@ static void print_resources(const Compiler &compiler, const char *tag, const Sma
compiler.get_decoration_bitset(type.self).get(DecorationBufferBlock);
bool is_sized_block = is_block && (compiler.get_storage_class(res.id) == StorageClassUniform ||
compiler.get_storage_class(res.id) == StorageClassUniformConstant);
uint32_t fallback_id = !is_push_constant && is_block ? res.base_type_id : res.id;
ID fallback_id = !is_push_constant && is_block ? ID(res.base_type_id) : ID(res.id);

uint32_t block_size = 0;
uint32_t runtime_array_stride = 0;
Expand All @@ -268,7 +268,7 @@ static void print_resources(const Compiler &compiler, const char *tag, const Sma
for (auto arr : type.array)
array = join("[", arr ? convert_to_string(arr) : "", "]") + array;

fprintf(stderr, " ID %03u : %s%s", res.id,
fprintf(stderr, " ID %03u : %s%s", uint32_t(res.id),
!res.name.empty() ? res.name.c_str() : compiler.get_fallback_name(fallback_id).c_str(), array.c_str());

if (mask.get(DecorationLocation))
Expand Down Expand Up @@ -442,7 +442,7 @@ static void print_spec_constants(const Compiler &compiler)
fprintf(stderr, "Specialization constants\n");
fprintf(stderr, "==================\n\n");
for (auto &c : spec_constants)
fprintf(stderr, "ID: %u, Spec ID: %u\n", c.id, c.constant_id);
fprintf(stderr, "ID: %u, Spec ID: %u\n", uint32_t(c.id), c.constant_id);
fprintf(stderr, "==================\n\n");
}

Expand Down Expand Up @@ -522,6 +522,7 @@ struct CLIArguments
bool vulkan_glsl_disable_ext_samplerless_texture_functions = false;
bool emit_line_directives = false;
SmallVector<uint32_t> msl_discrete_descriptor_sets;
SmallVector<pair<uint32_t, uint32_t>> msl_dynamic_buffers;
SmallVector<PLSArg> pls_in;
SmallVector<PLSArg> pls_out;
SmallVector<Remap> remaps;
Expand Down Expand Up @@ -600,6 +601,7 @@ static void print_help()
"\t[--msl-multiview]\n"
"\t[--msl-view-index-from-device-index]\n"
"\t[--msl-dispatch-base]\n"
"\t[--msl-dynamic-buffer <set index> <binding>]\n"
"\t[--hlsl]\n"
"\t[--reflect]\n"
"\t[--shader-model]\n"
Expand Down Expand Up @@ -764,6 +766,9 @@ static string compile_iteration(const CLIArguments &args, std::vector<uint32_t>
msl_comp->set_msl_options(msl_opts);
for (auto &v : args.msl_discrete_descriptor_sets)
msl_comp->add_discrete_descriptor_set(v);
uint32_t i = 0;
for (auto &v : args.msl_dynamic_buffers)
msl_comp->add_dynamic_buffer(v.first, v.second, i++);
}
else if (args.hlsl)
compiler.reset(new CompilerHLSL(move(spirv_parser.get_parsed_ir())));
Expand Down Expand Up @@ -1086,6 +1091,13 @@ static int main_inner(int argc, char *argv[])
cbs.add("--msl-view-index-from-device-index",
[&args](CLIParser &) { args.msl_view_index_from_device_index = true; });
cbs.add("--msl-dispatch-base", [&args](CLIParser &) { args.msl_dispatch_base = true; });
cbs.add("--msl-dynamic-buffer", [&args](CLIParser &parser) {
args.msl_argument_buffers = true;
// Make sure next_uint() is called in-order.
uint32_t desc_set = parser.next_uint();
uint32_t binding = parser.next_uint();
args.msl_dynamic_buffers.push_back(make_pair(desc_set, binding));
});
cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); });
cbs.add("--rename-entry-point", [&args](CLIParser &parser) {
auto old_name = parser.next_string();
Expand Down
Original file line number Diff line number Diff line change
@@ -0,0 +1,24 @@
RWByteAddressBuffer _9 : register(u6, space0);
globallycoherent RasterizerOrderedByteAddressBuffer _42 : register(u3, space0);
RasterizerOrderedByteAddressBuffer _52 : register(u4, space0);
RWTexture2D<unorm float4> img4 : register(u5, space0);
RasterizerOrderedTexture2D<unorm float4> img : register(u0, space0);
RasterizerOrderedTexture2D<unorm float4> img3 : register(u2, space0);
RasterizerOrderedTexture2D<uint> img2 : register(u1, space0);

void frag_main()
{
_9.Store(0, uint(0));
img4[int2(1, 1)] = float4(1.0f, 0.0f, 0.0f, 1.0f);
img[int2(0, 0)] = img3[int2(0, 0)];
uint _39;
InterlockedAdd(img2[int2(0, 0)], 1u, _39);
_42.Store(0, uint(int(_42.Load(0)) + 42));
uint _55;
_42.InterlockedAnd(4, _52.Load(0), _55);
}

void main()
{
frag_main();
}
Original file line number Diff line number Diff line change
@@ -0,0 +1,11 @@
#include <metal_stdlib>
#include <simd/simd.h>

using namespace metal;

kernel void main0(texture2d<float, access::write> uImage [[texture(0)]], texture2d<float> uImageRead [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
int2 _17 = int2(gl_GlobalInvocationID.xy);
uImage.write(uImageRead.read(uint2(_17)), uint2(_17));
}

Original file line number Diff line number Diff line change
@@ -0,0 +1,89 @@
#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;
}

Original file line number Diff line number Diff line change
Expand Up @@ -5,17 +5,6 @@

using namespace metal;

enum class spvSwizzle : uint
{
none = 0,
zero,
one,
red,
green,
blue,
alpha
};

template<typename T> struct spvRemoveReference { typedef T type; };
template<typename T> struct spvRemoveReference<thread T&> { typedef T type; };
template<typename T> struct spvRemoveReference<thread T&&> { typedef T type; };
Expand All @@ -28,6 +17,17 @@ template<typename T> inline constexpr thread T&& spvForward(thread typename spvR
return static_cast<thread T&&>(x);
}

enum class spvSwizzle : uint
{
none = 0,
zero,
one,
red,
green,
blue,
alpha
};

template<typename T>
inline T spvGetSwizzle(vec<T, 4> x, T c, spvSwizzle s)
{
Expand Down Expand Up @@ -65,66 +65,6 @@ inline T spvTextureSwizzle(T x, uint s)
return spvTextureSwizzle(vec<T, 4>(x, 0, 0, 1), s).x;
}

// Wrapper function that swizzles texture gathers.
template<typename T, typename Tex, typename... Ts>
inline vec<T, 4> spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) METAL_CONST_ARG(c)
{
if (sw)
{
switch (spvSwizzle((sw >> (uint(c) * 8)) & 0xFF))
{
case spvSwizzle::none:
break;
case spvSwizzle::zero:
return vec<T, 4>(0, 0, 0, 0);
case spvSwizzle::one:
return vec<T, 4>(1, 1, 1, 1);
case spvSwizzle::red:
return t.gather(s, spvForward<Ts>(params)..., component::x);
case spvSwizzle::green:
return t.gather(s, spvForward<Ts>(params)..., component::y);
case spvSwizzle::blue:
return t.gather(s, spvForward<Ts>(params)..., component::z);
case spvSwizzle::alpha:
return t.gather(s, spvForward<Ts>(params)..., component::w);
}
}
switch (c)
{
case component::x:
return t.gather(s, spvForward<Ts>(params)..., component::x);
case component::y:
return t.gather(s, spvForward<Ts>(params)..., component::y);
case component::z:
return t.gather(s, spvForward<Ts>(params)..., component::z);
case component::w:
return t.gather(s, spvForward<Ts>(params)..., component::w);
}
}

// Wrapper function that swizzles depth texture gathers.
template<typename T, typename Tex, typename... Ts>
inline vec<T, 4> spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw)
{
if (sw)
{
switch (spvSwizzle(sw & 0xFF))
{
case spvSwizzle::none:
case spvSwizzle::red:
break;
case spvSwizzle::zero:
case spvSwizzle::green:
case spvSwizzle::blue:
case spvSwizzle::alpha:
return vec<T, 4>(0, 0, 0, 0);
case spvSwizzle::one:
return vec<T, 4>(1, 1, 1, 1);
}
}
return t.gather_compare(s, spvForward<Ts>(params)...);
}

kernel void main0(constant uint* spvSwizzleConstants [[buffer(30)]], texture2d<float> foo [[texture(0)]], texture2d<float, access::write> bar [[texture(1)]], sampler fooSmplr [[sampler(0)]])
{
constant uint& fooSwzl = spvSwizzleConstants[0];
Expand Down
Loading

0 comments on commit 3dd517b

Please sign in to comment.