diff --git a/CMakeLists.txt b/CMakeLists.txt index 5f835e89..ac307a48 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -243,7 +243,7 @@ set(spirv-cross-util-sources ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross_util.hpp) set(spirv-cross-abi-major 0) -set(spirv-cross-abi-minor 58) +set(spirv-cross-abi-minor 59) set(spirv-cross-abi-patch 0) set(SPIRV_CROSS_VERSION ${spirv-cross-abi-major}.${spirv-cross-abi-minor}.${spirv-cross-abi-patch}) diff --git a/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp b/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..26730163 --- /dev/null +++ b/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,77 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +template +struct spvDescriptor +{ + T value; +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvDescriptor* ptr) : ptr(ptr) + { + } + const device T& operator [] (size_t i) const + { + return ptr[i].value; + } + const device spvDescriptor* ptr; +}; + +struct SSBO +{ + float4 v[1]; +}; + +struct UBO +{ + float4 v[1024]; +}; + +struct UBOs +{ + float4 v; +}; + +struct SSBOIn +{ + float4 v[1024]; +}; + +struct SSBOIns +{ + float4 v; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); + +struct spvDescriptorSetBuffer0 +{ + device SSBO* o [[id(0)]]; + constant UBO* v [[id(1)]]; + spvDescriptor vs [[id(2)]][1] /* unsized array hack */; +}; + +struct spvDescriptorSetBuffer1 +{ + device SSBOIn* w [[id(0)]]; + spvDescriptor ws [[id(1)]][1] /* unsized array hack */; +}; + +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + spvDescriptorArray vs {spvDescriptorSet0.vs}; + spvDescriptorArray ws {spvDescriptorSet1.ws}; + + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x]; + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v; + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x]; + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v; +} + diff --git a/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp b/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..2471c894 --- /dev/null +++ b/reference/opt/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,109 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +template +struct spvDescriptor +{ + T value; +}; + +template +struct spvBufferDescriptor +{ + T value; + int length; + const device T& operator -> () const device + { + return value; + } + const device T& operator * () const device + { + return value; + } +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvDescriptor* ptr) : ptr(ptr) + { + } + const device T& operator [] (size_t i) const + { + return ptr[i].value; + } + const device spvDescriptor* ptr; +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvBufferDescriptor* ptr) : ptr(ptr) + { + } + const device T* operator [] (size_t i) const + { + return ptr[i].value; + } + const int length(int i) const + { + return ptr[i].length; + } + const device spvBufferDescriptor* ptr; +}; + +struct SSBO +{ + float4 v[1]; +}; + +struct UBO +{ + float4 v[1024]; +}; + +struct UBOs +{ + float4 v; +}; + +struct SSBOIn +{ + float4 v[1024]; +}; + +struct SSBOIns +{ + float4 v; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); + +struct spvDescriptorSetBuffer0 +{ + device SSBO* o [[id(0)]]; + constant UBO* v [[id(1)]]; + spvDescriptor vs [[id(2)]][1] /* unsized array hack */; +}; + +struct spvDescriptorSetBuffer1 +{ + device SSBOIn* w [[id(0)]]; + spvBufferDescriptor ws [[id(1)]][1] /* unsized array hack */; +}; + +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + spvDescriptorArray vs {spvDescriptorSet0.vs}; + spvDescriptorArray ws {spvDescriptorSet1.ws}; + + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet0.v).v[gl_WorkGroupID.x]; + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v; + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = (*spvDescriptorSet1.w).v[gl_WorkGroupID.x]; + (*spvDescriptorSet0.o).v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v; +} + diff --git a/reference/opt/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp b/reference/opt/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..932cc19d --- /dev/null +++ b/reference/opt/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,55 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +template +struct spvDescriptor +{ + T value; +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvDescriptor* ptr) : ptr(ptr) + { + } + const device T& operator [] (size_t i) const + { + return ptr[i].value; + } + const device spvDescriptor* ptr; +}; + +struct SSBO +{ + float4 v[1]; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); + +struct spvDescriptorSetBuffer0 +{ + device SSBO* m_13 [[id(0)]]; + texture2d T [[id(1)]]; + spvDescriptor> Ts [[id(2)]][1] /* unsized array hack */; +}; + +struct spvDescriptorSetBuffer1 +{ + sampler S [[id(0)]]; + spvDescriptor Ss [[id(1)]][1] /* unsized array hack */; +}; + +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + spvDescriptorArray> Ts {spvDescriptorSet0.Ts}; + spvDescriptorArray Ss {spvDescriptorSet1.Ss}; + + (*spvDescriptorSet0.m_13).v[gl_GlobalInvocationID.x] = spvDescriptorSet0.T.sample(spvDescriptorSet1.S, float2(0.5), level(0.0)); + (*spvDescriptorSet0.m_13).v[gl_GlobalInvocationID.x] = Ts[gl_WorkGroupID.x].sample(Ss[gl_WorkGroupID.x], float2(0.5), level(0.0)); +} + diff --git a/reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp b/reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..fc3df7ff --- /dev/null +++ b/reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,83 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +template +struct spvDescriptor +{ + T value; +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvDescriptor* ptr) : ptr(ptr) + { + } + const device T& operator [] (size_t i) const + { + return ptr[i].value; + } + const device spvDescriptor* ptr; +}; + +struct SSBO +{ + float4 v[1]; +}; + +struct UBO +{ + float4 v[1024]; +}; + +struct UBOs +{ + float4 v; +}; + +struct SSBOIn +{ + float4 v[1024]; +}; + +struct SSBOIns +{ + float4 v; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); + +struct spvDescriptorSetBuffer0 +{ + device SSBO* o [[id(0)]]; + constant UBO* v [[id(1)]]; + spvDescriptor vs [[id(2)]][1] /* unsized array hack */; +}; + +struct spvDescriptorSetBuffer1 +{ + device SSBOIn* w [[id(0)]]; + spvDescriptor ws [[id(1)]][1] /* unsized array hack */; +}; + +static inline __attribute__((always_inline)) +void in_func(device SSBO& o, thread uint3& gl_GlobalInvocationID, constant UBO& v, thread uint3& gl_WorkGroupID, const spvDescriptorArray vs, device SSBOIn& w, const spvDescriptorArray ws) +{ + o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v; + o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v; +} + +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + spvDescriptorArray vs {spvDescriptorSet0.vs}; + spvDescriptorArray ws {spvDescriptorSet1.ws}; + + in_func((*spvDescriptorSet0.o), gl_GlobalInvocationID, (*spvDescriptorSet0.v), gl_WorkGroupID, vs, (*spvDescriptorSet1.w), ws); +} + diff --git a/reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp b/reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..db7c7e92 --- /dev/null +++ b/reference/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,115 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +template +struct spvDescriptor +{ + T value; +}; + +template +struct spvBufferDescriptor +{ + T value; + int length; + const device T& operator -> () const device + { + return value; + } + const device T& operator * () const device + { + return value; + } +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvDescriptor* ptr) : ptr(ptr) + { + } + const device T& operator [] (size_t i) const + { + return ptr[i].value; + } + const device spvDescriptor* ptr; +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvBufferDescriptor* ptr) : ptr(ptr) + { + } + const device T* operator [] (size_t i) const + { + return ptr[i].value; + } + const int length(int i) const + { + return ptr[i].length; + } + const device spvBufferDescriptor* ptr; +}; + +struct SSBO +{ + float4 v[1]; +}; + +struct UBO +{ + float4 v[1024]; +}; + +struct UBOs +{ + float4 v; +}; + +struct SSBOIn +{ + float4 v[1024]; +}; + +struct SSBOIns +{ + float4 v; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); + +struct spvDescriptorSetBuffer0 +{ + device SSBO* o [[id(0)]]; + constant UBO* v [[id(1)]]; + spvDescriptor vs [[id(2)]][1] /* unsized array hack */; +}; + +struct spvDescriptorSetBuffer1 +{ + device SSBOIn* w [[id(0)]]; + spvBufferDescriptor ws [[id(1)]][1] /* unsized array hack */; +}; + +static inline __attribute__((always_inline)) +void in_func(device SSBO& o, thread uint3& gl_GlobalInvocationID, constant UBO& v, thread uint3& gl_WorkGroupID, const spvDescriptorArray vs, device SSBOIn& w, const spvDescriptorArray ws) +{ + o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x]->v; + o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x]->v; +} + +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + spvDescriptorArray vs {spvDescriptorSet0.vs}; + spvDescriptorArray ws {spvDescriptorSet1.ws}; + + in_func((*spvDescriptorSet0.o), gl_GlobalInvocationID, (*spvDescriptorSet0.v), gl_WorkGroupID, vs, (*spvDescriptorSet1.w), ws); +} + diff --git a/reference/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp b/reference/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..7ff6cd77 --- /dev/null +++ b/reference/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,61 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +template +struct spvDescriptor +{ + T value; +}; + +template +struct spvDescriptorArray +{ + spvDescriptorArray(const device spvDescriptor* ptr) : ptr(ptr) + { + } + const device T& operator [] (size_t i) const + { + return ptr[i].value; + } + const device spvDescriptor* ptr; +}; + +struct SSBO +{ + float4 v[1]; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u); + +struct spvDescriptorSetBuffer0 +{ + device SSBO* m_13 [[id(0)]]; + texture2d T [[id(1)]]; + spvDescriptor> Ts [[id(2)]][1] /* unsized array hack */; +}; + +struct spvDescriptorSetBuffer1 +{ + sampler S [[id(0)]]; + spvDescriptor Ss [[id(1)]][1] /* unsized array hack */; +}; + +static inline __attribute__((always_inline)) +void in_func(device SSBO& _13, thread uint3& gl_GlobalInvocationID, texture2d T, sampler S, const spvDescriptorArray> Ts, thread uint3& gl_WorkGroupID, const spvDescriptorArray Ss) +{ + _13.v[gl_GlobalInvocationID.x] = T.sample(S, float2(0.5), level(0.0)); + _13.v[gl_GlobalInvocationID.x] = Ts[gl_WorkGroupID.x].sample(Ss[gl_WorkGroupID.x], float2(0.5), level(0.0)); +} + +kernel void main0(const device spvDescriptorSetBuffer0& spvDescriptorSet0 [[buffer(0)]], const device spvDescriptorSetBuffer1& spvDescriptorSet1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]]) +{ + spvDescriptorArray> Ts {spvDescriptorSet0.Ts}; + spvDescriptorArray Ss {spvDescriptorSet1.Ss}; + + in_func((*spvDescriptorSet0.m_13), gl_GlobalInvocationID, spvDescriptorSet0.T, spvDescriptorSet1.S, Ts, gl_WorkGroupID, Ss); +} + diff --git a/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp b/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..6f0d06a8 --- /dev/null +++ b/shaders-msl/comp/argument-buffers-runtime-array-buffer.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,28 @@ +#version 450 +#extension GL_EXT_nonuniform_qualifier : require +layout(local_size_x = 64) in; + +layout(set = 0, binding = 0) buffer SSBO +{ + vec4 v[]; +} o; + +layout(set = 0, binding = 1) uniform UBO { vec4 v[1024]; } v; +layout(set = 0, binding = 2) uniform UBOs { vec4 v; } vs[]; + +layout(set = 1, binding = 0) buffer SSBOIn { vec4 v[1024]; } w; +layout(set = 1, binding = 1) buffer SSBOIns { vec4 v; } ws[]; + +void in_func() +{ + o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x].v; + + o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x].v; +} + +void main() +{ + in_func(); +} diff --git a/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp b/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..6f0d06a8 --- /dev/null +++ b/shaders-msl/comp/argument-buffers-runtime-array-buffer.rich-descriptor.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,28 @@ +#version 450 +#extension GL_EXT_nonuniform_qualifier : require +layout(local_size_x = 64) in; + +layout(set = 0, binding = 0) buffer SSBO +{ + vec4 v[]; +} o; + +layout(set = 0, binding = 1) uniform UBO { vec4 v[1024]; } v; +layout(set = 0, binding = 2) uniform UBOs { vec4 v; } vs[]; + +layout(set = 1, binding = 0) buffer SSBOIn { vec4 v[1024]; } w; +layout(set = 1, binding = 1) buffer SSBOIns { vec4 v; } ws[]; + +void in_func() +{ + o.v[gl_GlobalInvocationID.x] = v.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = vs[gl_WorkGroupID.x].v; + + o.v[gl_GlobalInvocationID.x] = w.v[gl_WorkGroupID.x]; + o.v[gl_GlobalInvocationID.x] = ws[gl_WorkGroupID.x].v; +} + +void main() +{ + in_func(); +} diff --git a/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp b/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp new file mode 100644 index 00000000..58144ffa --- /dev/null +++ b/shaders-msl/comp/argument-buffers-runtime-array.argument.device-argument-buffer.argument-tier-1.msl2.comp @@ -0,0 +1,25 @@ +#version 450 +#extension GL_EXT_nonuniform_qualifier : require +layout(local_size_x = 64) in; + +layout(set = 0, binding = 0) buffer SSBO +{ + vec4 v[]; +}; + +layout(set = 0, binding = 1) uniform texture2D T; +layout(set = 0, binding = 2) uniform texture2D Ts[]; + +layout(set = 1, binding = 0) uniform sampler S; +layout(set = 1, binding = 1) uniform sampler Ss[]; + +void in_func() +{ + v[gl_GlobalInvocationID.x] = textureLod(sampler2D(T, S), vec2(0.5), 0.0); + v[gl_GlobalInvocationID.x] = textureLod(sampler2D(Ts[gl_WorkGroupID.x], Ss[gl_WorkGroupID.x]), vec2(0.5), 0.0); +} + +void main() +{ + in_func(); +} diff --git a/spirv_common.hpp b/spirv_common.hpp index 7fce7ae7..0cdbda23 100644 --- a/spirv_common.hpp +++ b/spirv_common.hpp @@ -1696,6 +1696,7 @@ struct Meta uint32_t index = 0; spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax; bool builtin = false; + bool qualified_alias_explicit_override = false; struct Extended { diff --git a/spirv_cross_c.cpp b/spirv_cross_c.cpp index dc2e5ea8..679333da 100644 --- a/spirv_cross_c.cpp +++ b/spirv_cross_c.cpp @@ -2552,6 +2552,16 @@ int spvc_constant_get_scalar_i32(spvc_constant constant, unsigned column, unsign return constant->scalar_i32(column, row); } +unsigned long long spvc_constant_get_scalar_u64(spvc_constant constant, unsigned column, unsigned row) +{ + return constant->scalar_u64(column, row); +} + +long long spvc_constant_get_scalar_i64(spvc_constant constant, unsigned column, unsigned row) +{ + return constant->scalar_i64(column, row); +} + unsigned spvc_constant_get_scalar_u16(spvc_constant constant, unsigned column, unsigned row) { return constant->scalar_u16(column, row); @@ -2609,6 +2619,16 @@ void spvc_constant_set_scalar_i32(spvc_constant constant, unsigned column, unsig constant->m.c[column].r[row].i32 = value; } +void spvc_constant_set_scalar_u64(spvc_constant constant, unsigned column, unsigned row, unsigned long long value) +{ + constant->m.c[column].r[row].u64 = value; +} + +void spvc_constant_set_scalar_i64(spvc_constant constant, unsigned column, unsigned row, long long value) +{ + constant->m.c[column].r[row].i64 = value; +} + void spvc_constant_set_scalar_u16(spvc_constant constant, unsigned column, unsigned row, unsigned short value) { constant->m.c[column].r[row].u32 = uint32_t(value); diff --git a/spirv_cross_c.h b/spirv_cross_c.h index dd9884f5..47369964 100644 --- a/spirv_cross_c.h +++ b/spirv_cross_c.h @@ -40,7 +40,7 @@ extern "C" { /* Bumped if ABI or API breaks backwards compatibility. */ #define SPVC_C_API_VERSION_MAJOR 0 /* Bumped if APIs or enumerations are added in a backwards compatible way. */ -#define SPVC_C_API_VERSION_MINOR 58 +#define SPVC_C_API_VERSION_MINOR 59 /* Bumped if internal implementation details change. */ #define SPVC_C_API_VERSION_PATCH 0 @@ -1052,6 +1052,8 @@ SPVC_PUBLIC_API int spvc_constant_get_scalar_i16(spvc_constant constant, unsigne SPVC_PUBLIC_API unsigned spvc_constant_get_scalar_u8(spvc_constant constant, unsigned column, unsigned row); SPVC_PUBLIC_API int spvc_constant_get_scalar_i8(spvc_constant constant, unsigned column, unsigned row); SPVC_PUBLIC_API void spvc_constant_get_subconstants(spvc_constant constant, const spvc_constant_id **constituents, size_t *count); +SPVC_PUBLIC_API unsigned long long spvc_constant_get_scalar_u64(spvc_constant constant, unsigned column, unsigned row); +SPVC_PUBLIC_API long long spvc_constant_get_scalar_i64(spvc_constant constant, unsigned column, unsigned row); SPVC_PUBLIC_API spvc_type_id spvc_constant_get_type(spvc_constant constant); /* @@ -1062,6 +1064,8 @@ SPVC_PUBLIC_API void spvc_constant_set_scalar_fp32(spvc_constant constant, unsig SPVC_PUBLIC_API void spvc_constant_set_scalar_fp64(spvc_constant constant, unsigned column, unsigned row, double value); SPVC_PUBLIC_API void spvc_constant_set_scalar_u32(spvc_constant constant, unsigned column, unsigned row, unsigned value); SPVC_PUBLIC_API void spvc_constant_set_scalar_i32(spvc_constant constant, unsigned column, unsigned row, int value); +SPVC_PUBLIC_API void spvc_constant_set_scalar_u64(spvc_constant constant, unsigned column, unsigned row, unsigned long long value); +SPVC_PUBLIC_API void spvc_constant_set_scalar_i64(spvc_constant constant, unsigned column, unsigned row, long long value); SPVC_PUBLIC_API void spvc_constant_set_scalar_u16(spvc_constant constant, unsigned column, unsigned row, unsigned short value); SPVC_PUBLIC_API void spvc_constant_set_scalar_i16(spvc_constant constant, unsigned column, unsigned row, signed short value); SPVC_PUBLIC_API void spvc_constant_set_scalar_u8(spvc_constant constant, unsigned column, unsigned row, unsigned char value); diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index a7d2510f..08884486 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -15543,7 +15543,7 @@ string CompilerGLSL::type_to_array_glsl(const SPIRType &type) } } -string CompilerGLSL::image_type_glsl(const SPIRType &type, uint32_t id) +string CompilerGLSL::image_type_glsl(const SPIRType &type, uint32_t id, bool /*member*/) { auto &imagetype = get(type.image.type); string res; diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index c04ef88b..4cba0bb6 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -438,7 +438,7 @@ protected: virtual void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index, const std::string &qualifier = "", uint32_t base_offset = 0); virtual void emit_struct_padding_target(const SPIRType &type); - virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0); + virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0, bool member = false); std::string constant_expression(const SPIRConstant &c, bool inside_block_like_struct_scope = false, bool inside_struct_scope = false); diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 200167a2..c64484c9 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -1379,20 +1379,28 @@ void CompilerMSL::emit_entry_point_declarations() { SPIRV_CROSS_THROW("Unsized array of descriptors requires argument buffer tier 2"); } + + string resource_name; + if (descriptor_set_is_argument_buffer(get_decoration(var.self, DecorationDescriptorSet))) + resource_name = ir.meta[var.self].decoration.qualified_alias; + else + resource_name = name + "_"; + switch (type.basetype) { case SPIRType::Image: case SPIRType::Sampler: case SPIRType::AccelerationStructure: - statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", name, "_};"); + statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", resource_name, "};"); break; case SPIRType::SampledImage: - statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", name, "_};"); + statement("spvDescriptorArray<", type_to_glsl(buffer_type), "> ", name, " {", resource_name, "};"); + // Unsupported with argument buffer for now. statement("spvDescriptorArray ", name, "Smplr {", name, "Smplr_};"); break; case SPIRType::Struct: statement("spvDescriptorArray<", get_argument_address_space(var), " ", type_to_glsl(buffer_type), "*> ", - name, " {", name, "_};"); + name, " {", resource_name, "};"); break; default: break; @@ -10665,7 +10673,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) { if (arg_type.array.empty() || (var ? is_var_runtime_size_array(*var) : is_runtime_size_array(arg_type))) { - decl += join(", ", sampler_type(arg_type, arg.id), " ", to_sampler_expression(name_id)); + decl += join(", ", sampler_type(arg_type, arg.id, false), " ", to_sampler_expression(name_id)); } else { @@ -10673,7 +10681,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) descriptor_address_space(name_id, StorageClassUniformConstant, "thread const"); - decl += join(", ", sampler_address_space, " ", sampler_type(arg_type, name_id), "& ", + decl += join(", ", sampler_address_space, " ", sampler_type(arg_type, name_id, false), "& ", to_sampler_expression(name_id)); } } @@ -12063,6 +12071,20 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ array_type = type_to_array_glsl(physical_type); } + if (orig_id) + { + auto *data_type = declared_type; + if (is_pointer(*data_type)) + data_type = &get_pointee_type(*data_type); + + if (is_array(*data_type) && get_resource_array_size(*data_type, orig_id) == 0) + { + // Hack for declaring unsized array of resources. Need to declare dummy sized array by value inline. + // This can then be wrapped in spvDescriptorArray as usual. + array_type = "[1] /* unsized array hack */"; + } + } + string decl_type; if (declared_type->vecsize > 4) { @@ -13283,12 +13305,22 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) !is_hidden_variable(var)) { auto &type = get_variable_data_type(var); + uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); if (is_supported_argument_buffer_type(type) && var.storage != StorageClassPushConstant) { - uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); if (descriptor_set_is_argument_buffer(desc_set)) + { + if (is_var_runtime_size_array(var)) + { + // Runtime arrays need to be wrapped in spvDescriptorArray from argument buffer payload. + entry_point_bindings.push_back(&var); + // We'll wrap this, so to_name() will always use non-qualified name. + // We'll need the qualified name to create temporary variable instead. + ir.meta[var_id].decoration.qualified_alias_explicit_override = true; + } return; + } } // Handle descriptor aliasing. We can handle aliasing of buffers by casting pointers, @@ -13480,7 +13512,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) case SPIRType::Sampler: if (!ep_args.empty()) ep_args += ", "; - ep_args += sampler_type(type, var_id) + " " + r.name; + ep_args += sampler_type(type, var_id, false) + " " + r.name; if (is_var_runtime_size_array(var)) ep_args += "_ [[buffer(" + convert_to_string(r.index) + ")]]"; else @@ -13495,7 +13527,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) const auto &basetype = get(var.basetype); if (!type_is_msl_framebuffer_fetch(basetype)) { - ep_args += image_type_glsl(type, var_id) + " " + r.name; + ep_args += image_type_glsl(type, var_id, false) + " " + r.name; if (r.plane > 0) ep_args += join(plane_name_suffix, r.plane); @@ -13512,7 +13544,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) { if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 3)) SPIRV_CROSS_THROW("Framebuffer fetch on Mac is not supported before MSL 2.3."); - ep_args += image_type_glsl(type, var_id) + " " + r.name; + ep_args += image_type_glsl(type, var_id, false) + " " + r.name; ep_args += " [[color(" + convert_to_string(r.index) + ")]]"; } @@ -14500,11 +14532,15 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg) if (!address_space.empty()) decl = join(address_space, " ", decl); - const char *argument_buffer_space = descriptor_address_space(name_id, type_storage, nullptr); - if (argument_buffer_space) + // spvDescriptorArray absorbs the address space inside the template. + if (!is_var_runtime_size_array(var)) { - decl += " "; - decl += argument_buffer_space; + const char *argument_buffer_space = descriptor_address_space(name_id, type_storage, nullptr); + if (argument_buffer_space) + { + decl += " "; + decl += argument_buffer_space; + } } // Special case, need to override the array size here if we're using tess level as an argument. @@ -14626,7 +14662,7 @@ string CompilerMSL::to_name(uint32_t id, bool allow_alias) const if (current_function && (current_function->self == ir.default_entry_point)) { auto *m = ir.find_meta(id); - if (m && !m->decoration.qualified_alias.empty()) + if (m && !m->decoration.qualified_alias_explicit_override && !m->decoration.qualified_alias.empty()) return m->decoration.qualified_alias; } return Compiler::to_name(id, allow_alias); @@ -15049,6 +15085,26 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member) auto type_address_space = get_type_address_space(type, id); const auto *p_parent_type = &get(type.parent_type); + // If we're wrapping buffer descriptors in a spvDescriptorArray, we'll have to handle it as a special case. + if (member && id) + { + auto &var = get(id); + if (is_var_runtime_size_array(var) && is_runtime_size_array(*p_parent_type)) + { + const bool ssbo = has_decoration(p_parent_type->self, DecorationBufferBlock); + bool buffer_desc = + (var.storage == StorageClassStorageBuffer || ssbo) && + msl_options.runtime_array_rich_descriptor; + + const char *wrapper_type = buffer_desc ? "spvBufferDescriptor" : "spvDescriptor"; + add_spv_func_and_recompile(SPVFuncImplVariableDescriptorArray); + add_spv_func_and_recompile(buffer_desc ? SPVFuncImplVariableSizedDescriptor : SPVFuncImplVariableDescriptor); + + type_name = join(wrapper_type, "<", type_address_space, " ", type_to_glsl(*p_parent_type, id), " *>"); + return type_name; + } + } + // Work around C pointer qualifier rules. If glsl_type is a pointer type as well // we'll need to emit the address space to the right. // We could always go this route, but it makes the code unnatural. @@ -15107,10 +15163,10 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member) case SPIRType::Image: case SPIRType::SampledImage: - return image_type_glsl(type, id); + return image_type_glsl(type, id, member); case SPIRType::Sampler: - return sampler_type(type, id); + return sampler_type(type, id, member); case SPIRType::Void: return "void"; @@ -15326,7 +15382,7 @@ std::string CompilerMSL::variable_decl(const SPIRType &type, const std::string & return CompilerGLSL::variable_decl(type, name, id); } -std::string CompilerMSL::sampler_type(const SPIRType &type, uint32_t id) +std::string CompilerMSL::sampler_type(const SPIRType &type, uint32_t id, bool member) { auto *var = maybe_get(id); if (var && var->basevariable) @@ -15345,26 +15401,31 @@ std::string CompilerMSL::sampler_type(const SPIRType &type, uint32_t id) // Arrays of samplers in MSL must be declared with a special array syntax ala C++11 std::array. // If we have a runtime array, it could be a variable-count descriptor set binding. + auto &parent = get(get_pointee_type(type).parent_type); uint32_t array_size = get_resource_array_size(type, id); + if (array_size == 0) { add_spv_func_and_recompile(SPVFuncImplVariableDescriptor); add_spv_func_and_recompile(SPVFuncImplVariableDescriptorArray); - auto &parent = get(get_pointee_type(type).parent_type); - if (processing_entry_point) - return join("const device spvDescriptor<", sampler_type(parent, id), ">*"); - return join("const spvDescriptorArray<", sampler_type(parent, id), ">"); - } - auto &parent = get(get_pointee_type(type).parent_type); - return join("array<", sampler_type(parent, id), ", ", array_size, ">"); + const char *descriptor_wrapper = processing_entry_point ? "const device spvDescriptor" : "const spvDescriptorArray"; + if (member) + descriptor_wrapper = "spvDescriptor"; + return join(descriptor_wrapper, "<", sampler_type(parent, id, false), ">", + processing_entry_point ? "*" : ""); + } + else + { + return join("array<", sampler_type(parent, id, false), ", ", array_size, ">"); + } } else return "sampler"; } // Returns an MSL string describing the SPIR-V image type -string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id) +string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id, bool member) { auto *var = maybe_get(id); if (var && var->basevariable) @@ -15395,17 +15456,29 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id) // Arrays of images in MSL must be declared with a special array syntax ala C++11 std::array. // If we have a runtime array, it could be a variable-count descriptor set binding. + auto &parent = get(get_pointee_type(type).parent_type); uint32_t array_size = get_resource_array_size(type, id); + if (array_size == 0) { add_spv_func_and_recompile(SPVFuncImplVariableDescriptor); add_spv_func_and_recompile(SPVFuncImplVariableDescriptorArray); - auto &parent = get(get_pointee_type(type).parent_type); - return join("const device spvDescriptor<", image_type_glsl(parent, id), ">*"); + const char *descriptor_wrapper = processing_entry_point ? "const device spvDescriptor" : "const spvDescriptorArray"; + if (member) + { + descriptor_wrapper = "spvDescriptor"; + // This requires a specialized wrapper type that packs image and sampler side by side. + // It is possible in theory. + if (type.basetype == SPIRType::SampledImage) + SPIRV_CROSS_THROW("Argument buffer runtime array currently not supported for combined image sampler."); + } + return join(descriptor_wrapper, "<", image_type_glsl(parent, id, false), ">", + processing_entry_point ? "*" : ""); + } + else + { + return join("array<", image_type_glsl(parent, id, false), ", ", array_size, ">"); } - - auto &parent = get(get_pointee_type(type).parent_type); - return join("array<", image_type_glsl(parent, id), ", ", array_size, ">"); } string img_type_name; @@ -17842,6 +17915,9 @@ void CompilerMSL::analyze_argument_buffers() auto &var = *resource.var; auto &type = get_variable_data_type(var); + if (is_var_runtime_size_array(var) && (argument_buffer_device_storage_mask & (1u << desc_set)) == 0) + SPIRV_CROSS_THROW("Runtime sized variables must be in device storage argument buffers."); + // If needed, synthesize and add padding members. // member_index and next_arg_buff_index are incremented when padding members are added. if (msl_options.pad_argument_buffer_resources) diff --git a/spirv_msl.hpp b/spirv_msl.hpp index 45109d88..06744b87 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -862,8 +862,8 @@ protected: // GCC workaround of lambdas calling protected functions (for older GCC versions) std::string variable_decl(const SPIRType &type, const std::string &name, uint32_t id = 0) override; - std::string image_type_glsl(const SPIRType &type, uint32_t id = 0) override; - std::string sampler_type(const SPIRType &type, uint32_t id); + std::string image_type_glsl(const SPIRType &type, uint32_t id, bool member) override; + std::string sampler_type(const SPIRType &type, uint32_t id, bool member); std::string builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClass storage) override; std::string to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_t id) override; std::string to_name(uint32_t id, bool allow_alias = true) const override;