mirror of https://github.com/bkaradzic/bgfx
Updated spirv-cross.
This commit is contained in:
Бранимир Караџић
parent
33ececeb96
commit
518192046e
|
|
@ -18,6 +18,6 @@ uint2 spvTexelBufferCoord(uint tc)
|
|||
|
||||
kernel void main0(constant cb& _6 [[buffer(7)]], texture2d<float, access::write> _buffer [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
_buffer.write(_6.value, spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
|
||||
_buffer.write(float4(_6.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
|
||||
}
|
||||
|
||||
|
|
|
|||
10
3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp
vendored
Normal file
10
3rdparty/spirv-cross/reference/opt/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp
vendored
Normal file
|
|
@ -0,0 +1,10 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
kernel void main0(texture2d<float, access::read_write> TargetTexture [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
||||
{
|
||||
TargetTexture.write((TargetTexture.read(uint2(gl_WorkGroupID.xy)).xy + float2(1.0)).xyyy, uint2((gl_WorkGroupID.xy + uint2(1u))));
|
||||
}
|
||||
|
||||
48
3rdparty/spirv-cross/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp
vendored
Normal file
48
3rdparty/spirv-cross/reference/opt/shaders-msl/comp/storage-buffer-std140-vector-array.comp
vendored
Normal file
|
|
@ -0,0 +1,48 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct Sub
|
||||
{
|
||||
float4 f[2];
|
||||
float4 f2[2];
|
||||
float3 f3[2];
|
||||
float4 f4[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Sub sub[2];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
float _153[2];
|
||||
_153[0] = _27.sub[gl_WorkGroupID.x].f[0].x;
|
||||
_153[1] = _27.sub[gl_WorkGroupID.x].f[1].x;
|
||||
float2 _154[2];
|
||||
_154[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy;
|
||||
_154[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy;
|
||||
float3 _155[2];
|
||||
_155[0] = _27.sub[gl_WorkGroupID.x].f3[0];
|
||||
_155[1] = _27.sub[gl_WorkGroupID.x].f3[1];
|
||||
float4 _156[2];
|
||||
_156[0] = _27.sub[gl_WorkGroupID.x].f4[0];
|
||||
_156[1] = _27.sub[gl_WorkGroupID.x].f4[1];
|
||||
_153[gl_GlobalInvocationID.x] += 1.0;
|
||||
_154[gl_GlobalInvocationID.x] += float2(2.0);
|
||||
_155[gl_GlobalInvocationID.x] += float3(3.0);
|
||||
_156[gl_GlobalInvocationID.x] += float4(4.0);
|
||||
_27.sub[gl_WorkGroupID.x].f[0].x = _153[0];
|
||||
_27.sub[gl_WorkGroupID.x].f[1].x = _153[1];
|
||||
_27.sub[gl_WorkGroupID.x].f2[0].xy = _154[0];
|
||||
_27.sub[gl_WorkGroupID.x].f2[1].xy = _154[1];
|
||||
_27.sub[gl_WorkGroupID.x].f3[0] = _155[0];
|
||||
_27.sub[gl_WorkGroupID.x].f3[1] = _155[1];
|
||||
_27.sub[gl_WorkGroupID.x].f4[0] = _156[0];
|
||||
_27.sub[gl_WorkGroupID.x].f4[1] = _156[1];
|
||||
_27.sub[0].f[0].x += 5.0;
|
||||
_27.sub[0].f2[1].xy += float2(5.0);
|
||||
}
|
||||
|
||||
|
|
@ -69,7 +69,7 @@ struct SSBO1
|
|||
|
||||
struct S0_1
|
||||
{
|
||||
float2 a[1];
|
||||
float4 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
|
|
@ -115,21 +115,21 @@ struct SSBO0
|
|||
Content_1 content;
|
||||
Content_1 content1[2];
|
||||
Content_1 content2;
|
||||
float array[1];
|
||||
float4 array[1];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO0& ssbo_140 [[buffer(0)]], device SSBO1& ssbo_430 [[buffer(1)]])
|
||||
{
|
||||
Content_1 _60 = ssbo_140.content;
|
||||
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
|
||||
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0].xy;
|
||||
ssbo_430.content.m0s[0].b = _60.m0s[0].b;
|
||||
ssbo_430.content.m1s[0].a = _60.m1s[0].a;
|
||||
ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a);
|
||||
ssbo_430.content.m1s[0].b = _60.m1s[0].b;
|
||||
ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
|
||||
ssbo_430.content.m2s[0].b = _60.m2s[0].b;
|
||||
ssbo_430.content.m0.a[0] = _60.m0.a[0];
|
||||
ssbo_430.content.m0.a[0] = _60.m0.a[0].xy;
|
||||
ssbo_430.content.m0.b = _60.m0.b;
|
||||
ssbo_430.content.m1.a = _60.m1.a;
|
||||
ssbo_430.content.m1.a = float3(_60.m1.a);
|
||||
ssbo_430.content.m1.b = _60.m1.b;
|
||||
ssbo_430.content.m2.a[0] = _60.m2.a[0];
|
||||
ssbo_430.content.m2.b = _60.m2.b;
|
||||
|
|
|
|||
|
|
@ -18,6 +18,6 @@ uint2 spvTexelBufferCoord(uint tc)
|
|||
|
||||
kernel void main0(constant cb& _6 [[buffer(7)]], texture2d<float, access::write> _buffer [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
|
||||
{
|
||||
_buffer.write(_6.value, spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
|
||||
_buffer.write(float4(_6.value), spvTexelBufferCoord(((32u * gl_WorkGroupID.x) + gl_LocalInvocationIndex)));
|
||||
}
|
||||
|
||||
|
|
|
|||
21
3rdparty/spirv-cross/reference/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp
vendored
Normal file
21
3rdparty/spirv-cross/reference/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp
vendored
Normal file
|
|
@ -0,0 +1,21 @@
|
|||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
void _main(thread const uint3& id, thread texture2d<float, access::read_write> TargetTexture)
|
||||
{
|
||||
float2 loaded = TargetTexture.read(uint2(id.xy)).xy;
|
||||
float2 storeTemp = loaded + float2(1.0);
|
||||
TargetTexture.write(storeTemp.xyyy, uint2((id.xy + uint2(1u))));
|
||||
}
|
||||
|
||||
kernel void main0(texture2d<float, access::read_write> TargetTexture [[texture(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
||||
{
|
||||
uint3 id = gl_WorkGroupID;
|
||||
uint3 param = id;
|
||||
_main(param, TargetTexture);
|
||||
}
|
||||
|
||||
|
|
@ -32,7 +32,7 @@ float4 _main(thread const float4& pos, constant buf& v_11)
|
|||
{
|
||||
int _46 = int(pos.x) % 16;
|
||||
Foo_1 foo;
|
||||
foo.a = v_11.results[_46].a;
|
||||
foo.a = float3(v_11.results[_46].a);
|
||||
foo.b = v_11.results[_46].b;
|
||||
return float4(dot(foo.a, v_11.bar.xyz), foo.b, 0.0, 0.0);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -30,7 +30,7 @@ constant uint3 gl_WorkGroupSize = uint3(32u, 1u, 1u);
|
|||
kernel void main0(device Buffer0& _15 [[buffer(1)]], device Buffer1& _34 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
T1_1 v;
|
||||
v.a = _15.buf0[0].a;
|
||||
v.a = float3(_15.buf0[0].a);
|
||||
v.b = _15.buf0[0].b;
|
||||
float x = v.b;
|
||||
_34.buf1[gl_GlobalInvocationID.x] = x;
|
||||
|
|
|
|||
53
3rdparty/spirv-cross/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp
vendored
Normal file
53
3rdparty/spirv-cross/reference/shaders-msl/comp/storage-buffer-std140-vector-array.comp
vendored
Normal file
|
|
@ -0,0 +1,53 @@
|
|||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct Sub
|
||||
{
|
||||
float4 f[2];
|
||||
float4 f2[2];
|
||||
float3 f3[2];
|
||||
float4 f4[2];
|
||||
};
|
||||
|
||||
struct Sub_1
|
||||
{
|
||||
float f[2];
|
||||
float2 f2[2];
|
||||
float3 f3[2];
|
||||
float4 f4[2];
|
||||
};
|
||||
|
||||
struct SSBO
|
||||
{
|
||||
Sub sub[2];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO& _27 [[buffer(0)]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
Sub_1 foo;
|
||||
foo.f[0] = _27.sub[gl_WorkGroupID.x].f[0].x;
|
||||
foo.f[1] = _27.sub[gl_WorkGroupID.x].f[1].x;
|
||||
foo.f2[0] = _27.sub[gl_WorkGroupID.x].f2[0].xy;
|
||||
foo.f2[1] = _27.sub[gl_WorkGroupID.x].f2[1].xy;
|
||||
foo.f3[0] = _27.sub[gl_WorkGroupID.x].f3[0];
|
||||
foo.f3[1] = _27.sub[gl_WorkGroupID.x].f3[1];
|
||||
foo.f4[0] = _27.sub[gl_WorkGroupID.x].f4[0];
|
||||
foo.f4[1] = _27.sub[gl_WorkGroupID.x].f4[1];
|
||||
foo.f[gl_GlobalInvocationID.x] += 1.0;
|
||||
foo.f2[gl_GlobalInvocationID.x] += float2(2.0);
|
||||
foo.f3[gl_GlobalInvocationID.x] += float3(3.0);
|
||||
foo.f4[gl_GlobalInvocationID.x] += float4(4.0);
|
||||
_27.sub[gl_WorkGroupID.x].f[0].x = foo.f[0];
|
||||
_27.sub[gl_WorkGroupID.x].f[1].x = foo.f[1];
|
||||
_27.sub[gl_WorkGroupID.x].f2[0].xy = foo.f2[0];
|
||||
_27.sub[gl_WorkGroupID.x].f2[1].xy = foo.f2[1];
|
||||
_27.sub[gl_WorkGroupID.x].f3[0] = foo.f3[0];
|
||||
_27.sub[gl_WorkGroupID.x].f3[1] = foo.f3[1];
|
||||
_27.sub[gl_WorkGroupID.x].f4[0] = foo.f4[0];
|
||||
_27.sub[gl_WorkGroupID.x].f4[1] = foo.f4[1];
|
||||
_27.sub[0].f[0].x += 5.0;
|
||||
_27.sub[0].f2[1].xy += float2(5.0);
|
||||
}
|
||||
|
||||
|
|
@ -69,7 +69,7 @@ struct SSBO1
|
|||
|
||||
struct S0_1
|
||||
{
|
||||
float2 a[1];
|
||||
float4 a[1];
|
||||
float b;
|
||||
};
|
||||
|
||||
|
|
@ -115,21 +115,21 @@ struct SSBO0
|
|||
Content_1 content;
|
||||
Content_1 content1[2];
|
||||
Content_1 content2;
|
||||
float array[1];
|
||||
float4 array[1];
|
||||
};
|
||||
|
||||
kernel void main0(device SSBO0& ssbo_140 [[buffer(0)]], device SSBO1& ssbo_430 [[buffer(1)]])
|
||||
{
|
||||
Content_1 _60 = ssbo_140.content;
|
||||
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0];
|
||||
ssbo_430.content.m0s[0].a[0] = _60.m0s[0].a[0].xy;
|
||||
ssbo_430.content.m0s[0].b = _60.m0s[0].b;
|
||||
ssbo_430.content.m1s[0].a = _60.m1s[0].a;
|
||||
ssbo_430.content.m1s[0].a = float3(_60.m1s[0].a);
|
||||
ssbo_430.content.m1s[0].b = _60.m1s[0].b;
|
||||
ssbo_430.content.m2s[0].a[0] = _60.m2s[0].a[0];
|
||||
ssbo_430.content.m2s[0].b = _60.m2s[0].b;
|
||||
ssbo_430.content.m0.a[0] = _60.m0.a[0];
|
||||
ssbo_430.content.m0.a[0] = _60.m0.a[0].xy;
|
||||
ssbo_430.content.m0.b = _60.m0.b;
|
||||
ssbo_430.content.m1.a = _60.m1.a;
|
||||
ssbo_430.content.m1.a = float3(_60.m1.a);
|
||||
ssbo_430.content.m1.b = _60.m1.b;
|
||||
ssbo_430.content.m2.a[0] = _60.m2.a[0];
|
||||
ssbo_430.content.m2.b = _60.m2.b;
|
||||
|
|
|
|||
|
|
@ -35,7 +35,7 @@ struct main0_out
|
|||
float4 _main(thread const VertexOutput& IN, constant CB0& v_26)
|
||||
{
|
||||
TestStruct_1 st;
|
||||
st.position = v_26.CB0[1].position;
|
||||
st.position = float3(v_26.CB0[1].position);
|
||||
st.radius = v_26.CB0[1].radius;
|
||||
float4 col = float4(st.position, st.radius);
|
||||
return col;
|
||||
|
|
|
|||
|
|
@ -43,7 +43,7 @@ vertex main0_out main0(main0_in in [[stage_in]], constant UBO& _21 [[buffer(0)]]
|
|||
for (int i = 0; i < 4; i++)
|
||||
{
|
||||
Light_1 light;
|
||||
light.Position = _21.lights[i].Position;
|
||||
light.Position = float3(_21.lights[i].Position);
|
||||
light.Radius = _21.lights[i].Radius;
|
||||
light.Color = _21.lights[i].Color;
|
||||
float3 L = in.aVertex.xyz - light.Position;
|
||||
|
|
|
|||
75
3rdparty/spirv-cross/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp
vendored
Normal file
75
3rdparty/spirv-cross/shaders-msl/asm/comp/image-load-store-short-vector.asm.comp
vendored
Normal file
|
|
@ -0,0 +1,75 @@
|
|||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 7
|
||||
; Bound: 44
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability StorageImageExtendedFormats
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %id_1
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpSource HLSL 500
|
||||
OpName %main "main"
|
||||
OpName %_main_vu3_ "@main(vu3;"
|
||||
OpName %id "id"
|
||||
OpName %loaded "loaded"
|
||||
OpName %TargetTexture "TargetTexture"
|
||||
OpName %storeTemp "storeTemp"
|
||||
OpName %id_0 "id"
|
||||
OpName %id_1 "id"
|
||||
OpName %param "param"
|
||||
OpDecorate %TargetTexture DescriptorSet 0
|
||||
OpDecorate %TargetTexture Binding 0
|
||||
OpDecorate %id_1 BuiltIn WorkgroupId
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
|
||||
%9 = OpTypeFunction %void %_ptr_Function_v3uint
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%_ptr_Function_v2float = OpTypePointer Function %v2float
|
||||
%17 = OpTypeImage %float 2D 0 0 0 2 Rg32f
|
||||
%_ptr_UniformConstant_17 = OpTypePointer UniformConstant %17
|
||||
%TargetTexture = OpVariable %_ptr_UniformConstant_17 UniformConstant
|
||||
%v2uint = OpTypeVector %uint 2
|
||||
%float_1 = OpConstant %float 1
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%id_1 = OpVariable %_ptr_Input_v3uint Input
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%id_0 = OpVariable %_ptr_Function_v3uint Function
|
||||
%param = OpVariable %_ptr_Function_v3uint Function
|
||||
%40 = OpLoad %v3uint %id_1
|
||||
OpStore %id_0 %40
|
||||
%42 = OpLoad %v3uint %id_0
|
||||
OpStore %param %42
|
||||
%43 = OpFunctionCall %void %_main_vu3_ %param
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%_main_vu3_ = OpFunction %void None %9
|
||||
%id = OpFunctionParameter %_ptr_Function_v3uint
|
||||
%12 = OpLabel
|
||||
%loaded = OpVariable %_ptr_Function_v2float Function
|
||||
%storeTemp = OpVariable %_ptr_Function_v2float Function
|
||||
%20 = OpLoad %17 %TargetTexture
|
||||
%22 = OpLoad %v3uint %id
|
||||
%23 = OpVectorShuffle %v2uint %22 %22 0 1
|
||||
%24 = OpImageRead %v2float %20 %23
|
||||
OpStore %loaded %24
|
||||
%26 = OpLoad %v2float %loaded
|
||||
%28 = OpCompositeConstruct %v2float %float_1 %float_1
|
||||
%29 = OpFAdd %v2float %26 %28
|
||||
OpStore %storeTemp %29
|
||||
%30 = OpLoad %17 %TargetTexture
|
||||
%31 = OpLoad %v3uint %id
|
||||
%32 = OpVectorShuffle %v2uint %31 %31 0 1
|
||||
%34 = OpCompositeConstruct %v2uint %uint_1 %uint_1
|
||||
%35 = OpIAdd %v2uint %32 %34
|
||||
%36 = OpLoad %v2float %storeTemp
|
||||
OpImageWrite %30 %35 %36
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
30
3rdparty/spirv-cross/shaders-msl/comp/storage-buffer-std140-vector-array.comp
vendored
Normal file
30
3rdparty/spirv-cross/shaders-msl/comp/storage-buffer-std140-vector-array.comp
vendored
Normal file
|
|
@ -0,0 +1,30 @@
|
|||
#version 450
|
||||
layout(local_size_x = 1) in;
|
||||
|
||||
struct Sub
|
||||
{
|
||||
float f[2];
|
||||
vec2 f2[2];
|
||||
vec3 f3[2];
|
||||
vec4 f4[2];
|
||||
};
|
||||
|
||||
layout(std140, binding = 0) buffer SSBO
|
||||
{
|
||||
Sub sub[2];
|
||||
};
|
||||
|
||||
void main()
|
||||
{
|
||||
Sub foo = sub[gl_WorkGroupID.x];
|
||||
|
||||
foo.f[gl_GlobalInvocationID.x] += 1.0;
|
||||
foo.f2[gl_GlobalInvocationID.x] += 2.0;
|
||||
foo.f3[gl_GlobalInvocationID.x] += 3.0;
|
||||
foo.f4[gl_GlobalInvocationID.x] += 4.0;
|
||||
sub[gl_WorkGroupID.x] = foo;
|
||||
|
||||
sub[0].f[0] += 5.0;
|
||||
sub[0].f2[1] += 5.0;
|
||||
}
|
||||
|
||||
|
|
@ -1339,6 +1339,7 @@ T &variant_set(Variant &var, P &&... args)
|
|||
|
||||
struct AccessChainMeta
|
||||
{
|
||||
uint32_t storage_packed_type = 0;
|
||||
bool need_transpose = false;
|
||||
bool storage_is_packed = false;
|
||||
bool storage_is_invariant = false;
|
||||
|
|
@ -1365,6 +1366,12 @@ struct Meta
|
|||
uint32_t index = 0;
|
||||
spv::FPRoundingMode fp_rounding_mode = spv::FPRoundingModeMax;
|
||||
bool builtin = false;
|
||||
|
||||
struct
|
||||
{
|
||||
uint32_t packed_type = 0;
|
||||
bool packed = false;
|
||||
} extended;
|
||||
};
|
||||
|
||||
Decoration decoration;
|
||||
|
|
|
|||
|
|
@ -169,7 +169,7 @@ string Compiler::to_name(uint32_t id, bool allow_alias) const
|
|||
{
|
||||
// If the alias master has been specially packed, we will have emitted a clean variant as well,
|
||||
// so skip the name aliasing here.
|
||||
if (!has_decoration(type.type_alias, DecorationCPacked))
|
||||
if (!has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked))
|
||||
return to_name(type.type_alias);
|
||||
}
|
||||
}
|
||||
|
|
@ -515,7 +515,7 @@ bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *
|
|||
|
||||
bool Compiler::is_scalar(const SPIRType &type) const
|
||||
{
|
||||
return type.vecsize == 1 && type.columns == 1;
|
||||
return type.basetype != SPIRType::Struct && type.vecsize == 1 && type.columns == 1;
|
||||
}
|
||||
|
||||
bool Compiler::is_vector(const SPIRType &type) const
|
||||
|
|
@ -868,7 +868,7 @@ void Compiler::fixup_type_alias()
|
|||
for (auto alias_itr = begin(type_ids); alias_itr != end(type_ids); ++alias_itr)
|
||||
{
|
||||
auto &type = get<SPIRType>(*alias_itr);
|
||||
if (type.type_alias != 0 && !has_decoration(type.type_alias, DecorationCPacked))
|
||||
if (type.type_alias != 0 && !has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked))
|
||||
{
|
||||
// We will skip declaring this type, so make sure the type_alias type comes before.
|
||||
auto master_itr = find(begin(type_ids), end(type_ids), type.type_alias);
|
||||
|
|
@ -1165,6 +1165,153 @@ void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argum
|
|||
ir.set_decoration(id, decoration, argument);
|
||||
}
|
||||
|
||||
void Compiler::set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value)
|
||||
{
|
||||
auto &dec = ir.meta[id].decoration;
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
dec.extended.packed = true;
|
||||
break;
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
dec.extended.packed_type = value;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
void Compiler::set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration,
|
||||
uint32_t value)
|
||||
{
|
||||
ir.meta[type].members.resize(max(ir.meta[type].members.size(), size_t(index) + 1));
|
||||
auto &dec = ir.meta[type].members[index];
|
||||
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
dec.extended.packed = true;
|
||||
break;
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
dec.extended.packed_type = value;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
uint32_t Compiler::get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const
|
||||
{
|
||||
auto *m = ir.find_meta(id);
|
||||
if (!m)
|
||||
return 0;
|
||||
|
||||
auto &dec = m->decoration;
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
return uint32_t(dec.extended.packed);
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
return dec.extended.packed_type;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
uint32_t Compiler::get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const
|
||||
{
|
||||
auto *m = ir.find_meta(type);
|
||||
if (!m)
|
||||
return 0;
|
||||
|
||||
if (index >= m->members.size())
|
||||
return 0;
|
||||
|
||||
auto &dec = m->members[index];
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
return uint32_t(dec.extended.packed);
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
return dec.extended.packed_type;
|
||||
}
|
||||
|
||||
return 0;
|
||||
}
|
||||
|
||||
bool Compiler::has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const
|
||||
{
|
||||
auto *m = ir.find_meta(id);
|
||||
if (!m)
|
||||
return false;
|
||||
|
||||
auto &dec = m->decoration;
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
return dec.extended.packed;
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
return dec.extended.packed_type != 0;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
bool Compiler::has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const
|
||||
{
|
||||
auto *m = ir.find_meta(type);
|
||||
if (!m)
|
||||
return false;
|
||||
|
||||
if (index >= m->members.size())
|
||||
return false;
|
||||
|
||||
auto &dec = m->members[index];
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
return dec.extended.packed;
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
return dec.extended.packed_type != 0;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
||||
void Compiler::unset_extended_decoration(uint32_t id, ExtendedDecorations decoration)
|
||||
{
|
||||
auto &dec = ir.meta[id].decoration;
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
dec.extended.packed = false;
|
||||
break;
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
dec.extended.packed_type = 0;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
void Compiler::unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration)
|
||||
{
|
||||
ir.meta[type].members.resize(max(ir.meta[type].members.size(), size_t(index) + 1));
|
||||
auto &dec = ir.meta[type].members[index];
|
||||
|
||||
switch (decoration)
|
||||
{
|
||||
case SPIRVCrossDecorationPacked:
|
||||
dec.extended.packed = false;
|
||||
break;
|
||||
|
||||
case SPIRVCrossDecorationPackedType:
|
||||
dec.extended.packed_type = 0;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
StorageClass Compiler::get_storage_class(uint32_t id) const
|
||||
{
|
||||
return get<SPIRVariable>(id).storage;
|
||||
|
|
|
|||
|
|
@ -114,6 +114,12 @@ struct EntryPoint
|
|||
spv::ExecutionModel execution_model;
|
||||
};
|
||||
|
||||
enum ExtendedDecorations
|
||||
{
|
||||
SPIRVCrossDecorationPacked,
|
||||
SPIRVCrossDecorationPackedType
|
||||
};
|
||||
|
||||
class Compiler
|
||||
{
|
||||
public:
|
||||
|
|
@ -946,6 +952,17 @@ protected:
|
|||
|
||||
bool image_is_comparison(const SPIRType &type, uint32_t id) const;
|
||||
|
||||
void set_extended_decoration(uint32_t id, ExtendedDecorations decoration, uint32_t value = 0);
|
||||
uint32_t get_extended_decoration(uint32_t id, ExtendedDecorations decoration) const;
|
||||
bool has_extended_decoration(uint32_t id, ExtendedDecorations decoration) const;
|
||||
void unset_extended_decoration(uint32_t id, ExtendedDecorations decoration);
|
||||
|
||||
void set_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration,
|
||||
uint32_t value = 0);
|
||||
uint32_t get_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const;
|
||||
bool has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const;
|
||||
void unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration);
|
||||
|
||||
private:
|
||||
// Used only to implement the old deprecated get_entry_point() interface.
|
||||
const SPIREntryPoint &get_first_entry_point(const std::string &name) const;
|
||||
|
|
|
|||
|
|
@ -696,7 +696,7 @@ void CompilerGLSL::emit_struct(SPIRType &type)
|
|||
// Type-punning with these types is legal, which complicates things
|
||||
// when we are storing struct and array types in an SSBO for example.
|
||||
// If the type master is packed however, we can no longer assume that the struct declaration will be redundant.
|
||||
if (type.type_alias != 0 && !has_decoration(type.type_alias, DecorationCPacked))
|
||||
if (type.type_alias != 0 && !has_extended_decoration(type.type_alias, SPIRVCrossDecorationPacked))
|
||||
return;
|
||||
|
||||
add_resource_name(type.self);
|
||||
|
|
@ -810,9 +810,9 @@ string CompilerGLSL::layout_for_member(const SPIRType &type, uint32_t index)
|
|||
SPIRV_CROSS_THROW("Component decoration is not supported in ES targets.");
|
||||
}
|
||||
|
||||
// DecorationCPacked is set by layout_for_variable earlier to mark that we need to emit offset qualifiers.
|
||||
// SPIRVCrossDecorationPacked is set by layout_for_variable earlier to mark that we need to emit offset qualifiers.
|
||||
// This is only done selectively in GLSL as needed.
|
||||
if (has_decoration(type.self, DecorationCPacked) && dec.decoration_flags.get(DecorationOffset))
|
||||
if (has_extended_decoration(type.self, SPIRVCrossDecorationPacked) && dec.decoration_flags.get(DecorationOffset))
|
||||
attr.push_back(join("offset = ", dec.offset));
|
||||
|
||||
if (attr.empty())
|
||||
|
|
@ -1370,7 +1370,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var)
|
|||
// layout_for_variable() will be called before the actual buffer emit.
|
||||
// The alternative is a full pass before codegen where we deduce this decoration,
|
||||
// but then we are just doing the exact same work twice, and more complexity.
|
||||
set_decoration(type.self, DecorationCPacked);
|
||||
set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
|
||||
}
|
||||
else
|
||||
{
|
||||
|
|
@ -1398,7 +1398,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var)
|
|||
if (!options.es && !options.vulkan_semantics && options.version < 440)
|
||||
require_extension_internal("GL_ARB_enhanced_layouts");
|
||||
|
||||
set_decoration(type.self, DecorationCPacked);
|
||||
set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
|
||||
}
|
||||
else if (buffer_is_packing_standard(type, BufferPackingStd430EnhancedLayout))
|
||||
{
|
||||
|
|
@ -1409,7 +1409,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var)
|
|||
if (!options.es && !options.vulkan_semantics && options.version < 440)
|
||||
require_extension_internal("GL_ARB_enhanced_layouts");
|
||||
|
||||
set_decoration(type.self, DecorationCPacked);
|
||||
set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
|
||||
}
|
||||
else
|
||||
{
|
||||
|
|
@ -2384,7 +2384,7 @@ void CompilerGLSL::handle_invalid_expression(uint32_t id)
|
|||
// by wrapping the expression in a constructor of the appropriate type.
|
||||
// GLSL does not support packed formats, so simply return the expression.
|
||||
// Subclasses that do will override
|
||||
string CompilerGLSL::unpack_expression_type(string expr_str, const SPIRType &)
|
||||
string CompilerGLSL::unpack_expression_type(string expr_str, const SPIRType &, uint32_t)
|
||||
{
|
||||
return expr_str;
|
||||
}
|
||||
|
|
@ -2486,26 +2486,28 @@ string CompilerGLSL::to_enclosed_expression(uint32_t id, bool register_expressio
|
|||
return enclose_expression(to_expression(id, register_expression_read));
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_unpacked_expression(uint32_t id)
|
||||
string CompilerGLSL::to_unpacked_expression(uint32_t id, bool register_expression_read)
|
||||
{
|
||||
// If we need to transpose, it will also take care of unpacking rules.
|
||||
auto *e = maybe_get<SPIRExpression>(id);
|
||||
bool need_transpose = e && e->need_transpose;
|
||||
if (!need_transpose && has_decoration(id, DecorationCPacked))
|
||||
return unpack_expression_type(to_expression(id), expression_type(id));
|
||||
if (!need_transpose && has_extended_decoration(id, SPIRVCrossDecorationPacked))
|
||||
return unpack_expression_type(to_expression(id, register_expression_read), expression_type(id),
|
||||
get_extended_decoration(id, SPIRVCrossDecorationPackedType));
|
||||
else
|
||||
return to_expression(id);
|
||||
return to_expression(id, register_expression_read);
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_enclosed_unpacked_expression(uint32_t id)
|
||||
string CompilerGLSL::to_enclosed_unpacked_expression(uint32_t id, bool register_expression_read)
|
||||
{
|
||||
// If we need to transpose, it will also take care of unpacking rules.
|
||||
auto *e = maybe_get<SPIRExpression>(id);
|
||||
bool need_transpose = e && e->need_transpose;
|
||||
if (!need_transpose && has_decoration(id, DecorationCPacked))
|
||||
return unpack_expression_type(to_expression(id), expression_type(id));
|
||||
if (!need_transpose && has_extended_decoration(id, SPIRVCrossDecorationPacked))
|
||||
return unpack_expression_type(to_expression(id, register_expression_read), expression_type(id),
|
||||
get_extended_decoration(id, SPIRVCrossDecorationPackedType));
|
||||
else
|
||||
return to_enclosed_expression(id);
|
||||
return to_enclosed_expression(id, register_expression_read);
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_dereferenced_expression(uint32_t id, bool register_expression_read)
|
||||
|
|
@ -2517,28 +2519,28 @@ string CompilerGLSL::to_dereferenced_expression(uint32_t id, bool register_expre
|
|||
return to_expression(id, register_expression_read);
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_pointer_expression(uint32_t id)
|
||||
string CompilerGLSL::to_pointer_expression(uint32_t id, bool register_expression_read)
|
||||
{
|
||||
auto &type = expression_type(id);
|
||||
if (type.pointer && expression_is_lvalue(id) && !should_dereference(id))
|
||||
return address_of_expression(to_enclosed_expression(id));
|
||||
return address_of_expression(to_enclosed_expression(id, register_expression_read));
|
||||
else
|
||||
return to_expression(id);
|
||||
return to_unpacked_expression(id, register_expression_read);
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_enclosed_pointer_expression(uint32_t id)
|
||||
string CompilerGLSL::to_enclosed_pointer_expression(uint32_t id, bool register_expression_read)
|
||||
{
|
||||
auto &type = expression_type(id);
|
||||
if (type.pointer && expression_is_lvalue(id) && !should_dereference(id))
|
||||
return address_of_expression(to_enclosed_expression(id));
|
||||
return address_of_expression(to_enclosed_expression(id, register_expression_read));
|
||||
else
|
||||
return to_enclosed_expression(id);
|
||||
return to_enclosed_unpacked_expression(id, register_expression_read);
|
||||
}
|
||||
|
||||
string CompilerGLSL::to_extract_component_expression(uint32_t id, uint32_t index)
|
||||
{
|
||||
auto expr = to_enclosed_expression(id);
|
||||
if (has_decoration(id, DecorationCPacked))
|
||||
if (has_extended_decoration(id, SPIRVCrossDecorationPacked))
|
||||
return join(expr, "[", index, "]");
|
||||
else
|
||||
return join(expr, ".", index_to_swizzle(index));
|
||||
|
|
@ -2581,7 +2583,7 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read)
|
|||
return to_enclosed_expression(e.base_expression) + e.expression;
|
||||
else if (e.need_transpose)
|
||||
{
|
||||
bool is_packed = has_decoration(id, DecorationCPacked);
|
||||
bool is_packed = has_extended_decoration(id, SPIRVCrossDecorationPacked);
|
||||
return convert_row_major_matrix(e.expression, get<SPIRType>(e.expression_type), is_packed);
|
||||
}
|
||||
else
|
||||
|
|
@ -2776,8 +2778,8 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
|||
|
||||
case OpCompositeExtract:
|
||||
{
|
||||
auto expr =
|
||||
access_chain_internal(cop.arguments[0], &cop.arguments[1], uint32_t(cop.arguments.size() - 1), true, false);
|
||||
auto expr = access_chain_internal(cop.arguments[0], &cop.arguments[1], uint32_t(cop.arguments.size() - 1),
|
||||
ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr);
|
||||
return expr;
|
||||
}
|
||||
|
||||
|
|
@ -4210,6 +4212,10 @@ void CompilerGLSL::emit_texture_op(const Instruction &i)
|
|||
if (is_legacy() && image_is_comparison(imgtype, img))
|
||||
expr += ".r";
|
||||
|
||||
// Deals with reads from MSL. We might need to downconvert to fewer components.
|
||||
if (op == OpImageRead)
|
||||
expr = remap_swizzle(get<SPIRType>(result_type), 4, expr);
|
||||
|
||||
emit_op(result_type, id, expr, forward);
|
||||
for (auto &inherit : inherited_expressions)
|
||||
inherit_expression_dependencies(id, inherit);
|
||||
|
|
@ -5482,10 +5488,15 @@ const char *CompilerGLSL::index_to_swizzle(uint32_t index)
|
|||
}
|
||||
|
||||
string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count,
|
||||
bool index_is_literal, bool chain_only, bool ptr_chain,
|
||||
AccessChainMeta *meta, bool register_expression_read)
|
||||
AccessChainFlags flags, AccessChainMeta *meta)
|
||||
{
|
||||
string expr;
|
||||
|
||||
bool index_is_literal = (flags & ACCESS_CHAIN_INDEX_IS_LITERAL_BIT) != 0;
|
||||
bool chain_only = (flags & ACCESS_CHAIN_CHAIN_ONLY_BIT) != 0;
|
||||
bool ptr_chain = (flags & ACCESS_CHAIN_PTR_CHAIN_BIT) != 0;
|
||||
bool register_expression_read = (flags & ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT) == 0;
|
||||
|
||||
if (!chain_only)
|
||||
expr = to_enclosed_expression(base, register_expression_read);
|
||||
|
||||
|
|
@ -5496,7 +5507,8 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
|||
|
||||
bool access_chain_is_arrayed = expr.find_first_of('[') != string::npos;
|
||||
bool row_major_matrix_needs_conversion = is_non_native_row_major_matrix(base);
|
||||
bool is_packed = has_decoration(base, DecorationCPacked);
|
||||
bool is_packed = has_extended_decoration(base, SPIRVCrossDecorationPacked);
|
||||
uint32_t packed_type = get_extended_decoration(base, SPIRVCrossDecorationPackedType);
|
||||
bool is_invariant = has_decoration(base, DecorationInvariant);
|
||||
bool pending_array_enclose = false;
|
||||
bool dimension_flatten = false;
|
||||
|
|
@ -5672,6 +5684,11 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
|||
is_invariant = true;
|
||||
|
||||
is_packed = member_is_packed_type(*type, index);
|
||||
if (is_packed)
|
||||
packed_type = get_extended_member_decoration(type->self, index, SPIRVCrossDecorationPackedType);
|
||||
else
|
||||
packed_type = 0;
|
||||
|
||||
row_major_matrix_needs_conversion = member_is_non_native_row_major_matrix(*type, index);
|
||||
type = &get<SPIRType>(type->member_types[index]);
|
||||
}
|
||||
|
|
@ -5683,6 +5700,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
|||
expr = convert_row_major_matrix(expr, *type, is_packed);
|
||||
row_major_matrix_needs_conversion = false;
|
||||
is_packed = false;
|
||||
packed_type = 0;
|
||||
}
|
||||
|
||||
expr += "[";
|
||||
|
|
@ -5722,6 +5740,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
|||
}
|
||||
|
||||
is_packed = false;
|
||||
packed_type = 0;
|
||||
type_id = type->parent_type;
|
||||
type = &get<SPIRType>(type_id);
|
||||
}
|
||||
|
|
@ -5741,6 +5760,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
|||
meta->need_transpose = row_major_matrix_needs_conversion;
|
||||
meta->storage_is_packed = is_packed;
|
||||
meta->storage_is_invariant = is_invariant;
|
||||
meta->storage_packed_type = packed_type;
|
||||
}
|
||||
|
||||
return expr;
|
||||
|
|
@ -5772,7 +5792,11 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32
|
|||
}
|
||||
else if (flattened_structs.count(base) && count > 0)
|
||||
{
|
||||
auto chain = access_chain_internal(base, indices, count, false, true, ptr_chain, nullptr, false).substr(1);
|
||||
AccessChainFlags flags = ACCESS_CHAIN_CHAIN_ONLY_BIT | ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT;
|
||||
if (ptr_chain)
|
||||
flags |= ACCESS_CHAIN_PTR_CHAIN_BIT;
|
||||
|
||||
auto chain = access_chain_internal(base, indices, count, flags, nullptr).substr(1);
|
||||
if (meta)
|
||||
{
|
||||
meta->need_transpose = false;
|
||||
|
|
@ -5782,7 +5806,10 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32
|
|||
}
|
||||
else
|
||||
{
|
||||
return access_chain_internal(base, indices, count, false, false, ptr_chain, meta, false);
|
||||
AccessChainFlags flags = ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT;
|
||||
if (ptr_chain)
|
||||
flags |= ACCESS_CHAIN_PTR_CHAIN_BIT;
|
||||
return access_chain_internal(base, indices, count, flags, meta);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
@ -6603,6 +6630,30 @@ void CompilerGLSL::handle_store_to_invariant_variable(uint32_t store_id, uint32_
|
|||
disallow_forwarding_in_expression_chain(*expr);
|
||||
}
|
||||
|
||||
void CompilerGLSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression)
|
||||
{
|
||||
auto rhs = to_pointer_expression(rhs_expression);
|
||||
|
||||
// Statements to OpStore may be empty if it is a struct with zero members. Just forward the store to /dev/null.
|
||||
if (!rhs.empty())
|
||||
{
|
||||
handle_store_to_invariant_variable(lhs_expression, rhs_expression);
|
||||
|
||||
auto lhs = to_dereferenced_expression(lhs_expression);
|
||||
|
||||
// We might need to bitcast in order to store to a builtin.
|
||||
bitcast_to_builtin_store(lhs_expression, rhs, expression_type(rhs_expression));
|
||||
|
||||
// Tries to optimize assignments like "<lhs> = <lhs> op expr".
|
||||
// While this is purely cosmetic, this is important for legacy ESSL where loop
|
||||
// variable increments must be in either i++ or i += const-expr.
|
||||
// Without this, we end up with i = i + 1, which is correct GLSL, but not correct GLES 2.0.
|
||||
if (!optimize_read_modify_write(expression_type(rhs_expression), lhs, rhs))
|
||||
statement(lhs, " = ", rhs, ";");
|
||||
register_write(lhs_expression);
|
||||
}
|
||||
}
|
||||
|
||||
void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
{
|
||||
auto ops = stream(instruction);
|
||||
|
|
@ -6674,8 +6725,12 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
register_read(id, ptr, forward);
|
||||
|
||||
// Pass through whether the result is of a packed type.
|
||||
if (has_decoration(ptr, DecorationCPacked))
|
||||
set_decoration(id, DecorationCPacked);
|
||||
if (has_extended_decoration(ptr, SPIRVCrossDecorationPacked))
|
||||
{
|
||||
set_extended_decoration(id, SPIRVCrossDecorationPacked);
|
||||
set_extended_decoration(id, SPIRVCrossDecorationPackedType,
|
||||
get_extended_decoration(ptr, SPIRVCrossDecorationPackedType));
|
||||
}
|
||||
|
||||
inherit_expression_dependencies(id, ptr);
|
||||
if (forward)
|
||||
|
|
@ -6706,14 +6761,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
|
||||
// Mark the result as being packed. Some platforms handled packed vectors differently than non-packed.
|
||||
if (meta.storage_is_packed)
|
||||
set_decoration(ops[1], DecorationCPacked);
|
||||
else
|
||||
unset_decoration(ops[1], DecorationCPacked);
|
||||
|
||||
set_extended_decoration(ops[1], SPIRVCrossDecorationPacked);
|
||||
if (meta.storage_packed_type != 0)
|
||||
set_extended_decoration(ops[1], SPIRVCrossDecorationPackedType, meta.storage_packed_type);
|
||||
if (meta.storage_is_invariant)
|
||||
set_decoration(ops[1], DecorationInvariant);
|
||||
else
|
||||
unset_decoration(ops[1], DecorationInvariant);
|
||||
|
||||
for (uint32_t i = 2; i < length; i++)
|
||||
{
|
||||
|
|
@ -6743,27 +6795,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
}
|
||||
else
|
||||
{
|
||||
auto rhs = to_pointer_expression(ops[1]);
|
||||
|
||||
// Statements to OpStore may be empty if it is a struct with zero members. Just forward the store to /dev/null.
|
||||
if (!rhs.empty())
|
||||
{
|
||||
handle_store_to_invariant_variable(ops[0], ops[1]);
|
||||
|
||||
auto lhs = to_dereferenced_expression(ops[0]);
|
||||
|
||||
// We might need to bitcast in order to store to a builtin.
|
||||
bitcast_to_builtin_store(ops[0], rhs, expression_type(ops[1]));
|
||||
|
||||
// Tries to optimize assignments like "<lhs> = <lhs> op expr".
|
||||
// While this is purely cosmetic, this is important for legacy ESSL where loop
|
||||
// variable increments must be in either i++ or i += const-expr.
|
||||
// Without this, we end up with i = i + 1, which is correct GLSL, but not correct GLES 2.0.
|
||||
if (!optimize_read_modify_write(expression_type(ops[1]), lhs, rhs))
|
||||
statement(lhs, " = ", rhs, ";");
|
||||
register_write(ops[0]);
|
||||
}
|
||||
emit_store_statement(ops[0], ops[1]);
|
||||
}
|
||||
|
||||
// Storing a pointer results in a variable pointer, so we must conservatively assume
|
||||
// we can write through it.
|
||||
if (expression_type(ops[1]).pointer)
|
||||
|
|
@ -6775,7 +6809,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
{
|
||||
uint32_t result_type = ops[0];
|
||||
uint32_t id = ops[1];
|
||||
auto e = access_chain_internal(ops[2], &ops[3], length - 3, true);
|
||||
auto e = access_chain_internal(ops[2], &ops[3], length - 3, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr);
|
||||
set<SPIRExpression>(id, e + ".length()", result_type, true);
|
||||
break;
|
||||
}
|
||||
|
|
@ -7007,7 +7041,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
// Make a copy, then use access chain to store the variable.
|
||||
statement(declare_temporary(result_type, id), to_expression(vec), ";");
|
||||
set<SPIRExpression>(id, to_name(id), result_type, true);
|
||||
auto chain = access_chain_internal(id, &index, 1, false);
|
||||
auto chain = access_chain_internal(id, &index, 1, 0, nullptr);
|
||||
statement(chain, " = ", to_expression(comp), ";");
|
||||
break;
|
||||
}
|
||||
|
|
@ -7017,7 +7051,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
uint32_t result_type = ops[0];
|
||||
uint32_t id = ops[1];
|
||||
|
||||
auto expr = access_chain_internal(ops[2], &ops[3], 1, false);
|
||||
auto expr = access_chain_internal(ops[2], &ops[3], 1, 0, nullptr);
|
||||
emit_op(result_type, id, expr, should_forward(ops[2]));
|
||||
inherit_expression_dependencies(id, ops[2]);
|
||||
inherit_expression_dependencies(id, ops[3]);
|
||||
|
|
@ -7041,7 +7075,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
allow_base_expression = false;
|
||||
|
||||
// Packed expressions cannot be split up.
|
||||
if (has_decoration(ops[2], DecorationCPacked))
|
||||
if (has_extended_decoration(ops[2], SPIRVCrossDecorationPacked))
|
||||
allow_base_expression = false;
|
||||
|
||||
AccessChainMeta meta;
|
||||
|
|
@ -7062,14 +7096,15 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
//
|
||||
// Including the base will prevent this and would trigger multiple reads
|
||||
// from expression causing it to be forced to an actual temporary in GLSL.
|
||||
auto expr = access_chain_internal(ops[2], &ops[3], length, true, true, false, &meta);
|
||||
auto expr = access_chain_internal(ops[2], &ops[3], length,
|
||||
ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_CHAIN_ONLY_BIT, &meta);
|
||||
e = &emit_op(result_type, id, expr, true, !expression_is_forwarded(ops[2]));
|
||||
inherit_expression_dependencies(id, ops[2]);
|
||||
e->base_expression = ops[2];
|
||||
}
|
||||
else
|
||||
{
|
||||
auto expr = access_chain_internal(ops[2], &ops[3], length, true, false, false, &meta);
|
||||
auto expr = access_chain_internal(ops[2], &ops[3], length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta);
|
||||
e = &emit_op(result_type, id, expr, should_forward(ops[2]), !expression_is_forwarded(ops[2]));
|
||||
inherit_expression_dependencies(id, ops[2]);
|
||||
}
|
||||
|
|
@ -7079,7 +7114,9 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
// instead of loading everything through an access chain.
|
||||
e->need_transpose = meta.need_transpose;
|
||||
if (meta.storage_is_packed)
|
||||
set_decoration(id, DecorationCPacked);
|
||||
set_extended_decoration(id, SPIRVCrossDecorationPacked);
|
||||
if (meta.storage_packed_type != 0)
|
||||
set_extended_decoration(id, SPIRVCrossDecorationPackedType, meta.storage_packed_type);
|
||||
if (meta.storage_is_invariant)
|
||||
set_decoration(id, DecorationInvariant);
|
||||
|
||||
|
|
@ -7100,7 +7137,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
// Make a copy, then use access chain to store the variable.
|
||||
statement(declare_temporary(result_type, id), to_expression(composite), ";");
|
||||
set<SPIRExpression>(id, to_name(id), result_type, true);
|
||||
auto chain = access_chain_internal(id, elems, length, true);
|
||||
auto chain = access_chain_internal(id, elems, length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr);
|
||||
statement(chain, " = ", to_expression(obj), ";");
|
||||
|
||||
break;
|
||||
|
|
@ -7167,7 +7204,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
|||
shuffle = true;
|
||||
|
||||
// Cannot use swizzles with packed expressions, force shuffle path.
|
||||
if (!shuffle && has_decoration(vec0, DecorationCPacked))
|
||||
if (!shuffle && has_extended_decoration(vec0, SPIRVCrossDecorationPacked))
|
||||
shuffle = true;
|
||||
|
||||
string expr;
|
||||
|
|
@ -8796,7 +8833,7 @@ bool CompilerGLSL::member_is_non_native_row_major_matrix(const SPIRType &type, u
|
|||
// GLSL does not define packed data types, but certain subclasses do.
|
||||
bool CompilerGLSL::member_is_packed_type(const SPIRType &type, uint32_t index) const
|
||||
{
|
||||
return has_member_decoration(type.self, index, DecorationCPacked);
|
||||
return has_extended_member_decoration(type.self, index, SPIRVCrossDecorationPacked);
|
||||
}
|
||||
|
||||
// Wraps the expression string in a function call that converts the
|
||||
|
|
|
|||
|
|
@ -51,6 +51,15 @@ struct PlsRemap
|
|||
PlsFormat format;
|
||||
};
|
||||
|
||||
enum AccessChainFlagBits
|
||||
{
|
||||
ACCESS_CHAIN_INDEX_IS_LITERAL_BIT = 1 << 0,
|
||||
ACCESS_CHAIN_CHAIN_ONLY_BIT = 1 << 1,
|
||||
ACCESS_CHAIN_PTR_CHAIN_BIT = 1 << 2,
|
||||
ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT = 1 << 3
|
||||
};
|
||||
typedef uint32_t AccessChainFlags;
|
||||
|
||||
class CompilerGLSL : public Compiler
|
||||
{
|
||||
public:
|
||||
|
|
@ -261,7 +270,7 @@ protected:
|
|||
virtual void emit_buffer_block(const SPIRVariable &type);
|
||||
virtual void emit_push_constant_block(const SPIRVariable &var);
|
||||
virtual void emit_uniform(const SPIRVariable &var);
|
||||
virtual std::string unpack_expression_type(std::string expr_str, const SPIRType &type);
|
||||
virtual std::string unpack_expression_type(std::string expr_str, const SPIRType &type, uint32_t packed_type_id);
|
||||
|
||||
std::unique_ptr<std::ostringstream> buffer;
|
||||
|
||||
|
|
@ -447,9 +456,10 @@ protected:
|
|||
bool expression_is_forwarded(uint32_t id);
|
||||
SPIRExpression &emit_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs,
|
||||
bool suppress_usage_tracking = false);
|
||||
std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, bool index_is_literal,
|
||||
bool chain_only = false, bool ptr_chain = false, AccessChainMeta *meta = nullptr,
|
||||
bool register_expression_read = true);
|
||||
|
||||
std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, AccessChainFlags flags,
|
||||
AccessChainMeta *meta);
|
||||
|
||||
std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type,
|
||||
AccessChainMeta *meta = nullptr, bool ptr_chain = false);
|
||||
|
||||
|
|
@ -476,11 +486,11 @@ protected:
|
|||
void append_global_func_args(const SPIRFunction &func, uint32_t index, std::vector<std::string> &arglist);
|
||||
std::string to_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_enclosed_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_unpacked_expression(uint32_t id);
|
||||
std::string to_enclosed_unpacked_expression(uint32_t id);
|
||||
std::string to_unpacked_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_enclosed_unpacked_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_dereferenced_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_pointer_expression(uint32_t id);
|
||||
std::string to_enclosed_pointer_expression(uint32_t id);
|
||||
std::string to_pointer_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_enclosed_pointer_expression(uint32_t id, bool register_expression_read = true);
|
||||
std::string to_extract_component_expression(uint32_t id, uint32_t index);
|
||||
std::string enclose_expression(const std::string &expr);
|
||||
std::string dereference_expression(const std::string &expr);
|
||||
|
|
@ -624,6 +634,7 @@ protected:
|
|||
void disallow_forwarding_in_expression_chain(const SPIRExpression &expr);
|
||||
|
||||
bool expression_is_constant_null(uint32_t id) const;
|
||||
virtual void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression);
|
||||
|
||||
private:
|
||||
void init()
|
||||
|
|
|
|||
|
|
@ -1835,7 +1835,7 @@ void CompilerHLSL::emit_struct_member(const SPIRType &type, uint32_t member_type
|
|||
string packing_offset;
|
||||
bool is_push_constant = type.storage == StorageClassPushConstant;
|
||||
|
||||
if ((has_decoration(type.self, DecorationCPacked) || is_push_constant) &&
|
||||
if ((has_extended_decoration(type.self, SPIRVCrossDecorationPacked) || is_push_constant) &&
|
||||
has_member_decoration(type.self, index, DecorationOffset))
|
||||
{
|
||||
uint32_t offset = memb[index].offset - base_offset;
|
||||
|
|
@ -1870,7 +1870,7 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var)
|
|||
if (type.array.empty())
|
||||
{
|
||||
if (buffer_is_packing_standard(type, BufferPackingHLSLCbufferPackOffset))
|
||||
set_decoration(type.self, DecorationCPacked);
|
||||
set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
|
||||
else
|
||||
SPIRV_CROSS_THROW("cbuffer cannot be expressed with either HLSL packing layout or packoffset.");
|
||||
|
||||
|
|
@ -1952,7 +1952,7 @@ void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var)
|
|||
auto &type = get<SPIRType>(var.basetype);
|
||||
|
||||
if (buffer_is_packing_standard(type, BufferPackingHLSLCbufferPackOffset, layout.start, layout.end))
|
||||
set_decoration(type.self, DecorationCPacked);
|
||||
set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
|
||||
else
|
||||
SPIRV_CROSS_THROW(
|
||||
"root constant cbuffer cannot be expressed with either HLSL packing layout or packoffset.");
|
||||
|
|
|
|||
|
|
@ -745,7 +745,7 @@ void CompilerMSL::mark_packable_structs()
|
|||
}
|
||||
|
||||
// If the specified type is a struct, it and any nested structs
|
||||
// are marked as packable with the DecorationCPacked decoration,
|
||||
// are marked as packable with the SPIRVCrossDecorationPacked decoration,
|
||||
void CompilerMSL::mark_as_packable(SPIRType &type)
|
||||
{
|
||||
// If this is not the base type (eg. it's a pointer or array), tunnel down
|
||||
|
|
@ -757,7 +757,7 @@ void CompilerMSL::mark_as_packable(SPIRType &type)
|
|||
|
||||
if (type.basetype == SPIRType::Struct)
|
||||
{
|
||||
set_decoration(type.self, DecorationCPacked);
|
||||
set_extended_decoration(type.self, SPIRVCrossDecorationPacked);
|
||||
|
||||
// Recurse
|
||||
size_t mbr_cnt = type.member_types.size();
|
||||
|
|
@ -1479,6 +1479,7 @@ uint32_t CompilerMSL::ensure_correct_attribute_type(uint32_t type_id, uint32_t l
|
|||
auto &base_type = set<SPIRType>(base_type_id);
|
||||
base_type = type;
|
||||
base_type.basetype = type.basetype == SPIRType::Short ? SPIRType::UShort : SPIRType::UInt;
|
||||
base_type.pointer = false;
|
||||
|
||||
if (!type.pointer)
|
||||
return base_type_id;
|
||||
|
|
@ -1508,6 +1509,7 @@ uint32_t CompilerMSL::ensure_correct_attribute_type(uint32_t type_id, uint32_t l
|
|||
auto &base_type = set<SPIRType>(base_type_id);
|
||||
base_type = type;
|
||||
base_type.basetype = SPIRType::UInt;
|
||||
base_type.pointer = false;
|
||||
|
||||
if (!type.pointer)
|
||||
return base_type_id;
|
||||
|
|
@ -1551,7 +1553,11 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
|
|||
for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++)
|
||||
{
|
||||
if (is_member_packable(ib_type, mbr_idx))
|
||||
set_member_decoration(ib_type_id, mbr_idx, DecorationCPacked);
|
||||
{
|
||||
set_extended_member_decoration(ib_type_id, mbr_idx, SPIRVCrossDecorationPacked);
|
||||
set_extended_member_decoration(ib_type_id, mbr_idx, SPIRVCrossDecorationPackedType,
|
||||
ib_type.member_types[mbr_idx]);
|
||||
}
|
||||
|
||||
// Align current offset to the current member's default alignment.
|
||||
size_t align_mask = get_declared_struct_member_alignment(ib_type, mbr_idx) - 1;
|
||||
|
|
@ -1580,11 +1586,30 @@ void CompilerMSL::align_struct(SPIRType &ib_type)
|
|||
bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
|
||||
{
|
||||
// We've already marked it as packable
|
||||
if (has_member_decoration(ib_type.self, index, DecorationCPacked))
|
||||
if (has_extended_member_decoration(ib_type.self, index, SPIRVCrossDecorationPacked))
|
||||
return true;
|
||||
|
||||
auto &mbr_type = get<SPIRType>(ib_type.member_types[index]);
|
||||
|
||||
uint32_t component_size = mbr_type.width / 8;
|
||||
uint32_t unpacked_mbr_size;
|
||||
if (mbr_type.vecsize == 3)
|
||||
unpacked_mbr_size = component_size * (mbr_type.vecsize + 1) * mbr_type.columns;
|
||||
else
|
||||
unpacked_mbr_size = component_size * mbr_type.vecsize * mbr_type.columns;
|
||||
|
||||
// Special case for packing. Check for float[] or vec2[] in std140 layout. Here we actually need to pad out instead,
|
||||
// but we will use the same mechanism.
|
||||
if (is_array(mbr_type) && (is_scalar(mbr_type) || is_vector(mbr_type)) && mbr_type.vecsize <= 2 &&
|
||||
type_struct_member_array_stride(ib_type, index) == 4 * component_size)
|
||||
{
|
||||
return true;
|
||||
}
|
||||
|
||||
// TODO: Another sanity check for matrices. We currently do not support std140 matrices which need to be padded out per column.
|
||||
//if (is_matrix(mbr_type) && mbr_type.vecsize <= 2 && type_struct_member_matrix_stride(ib_type, index) == 16)
|
||||
// SPIRV_CROSS_THROW("Currently cannot support matrices with small vector size in std140 layout.");
|
||||
|
||||
// Only vectors or 3-row matrices need to be packed.
|
||||
if (mbr_type.vecsize == 1 || (is_matrix(mbr_type) && mbr_type.vecsize != 3))
|
||||
return false;
|
||||
|
|
@ -1593,12 +1618,6 @@ bool CompilerMSL::is_member_packable(SPIRType &ib_type, uint32_t index)
|
|||
if (is_matrix(mbr_type) && !has_member_decoration(ib_type.self, index, DecorationRowMajor))
|
||||
return false;
|
||||
|
||||
uint32_t component_size = mbr_type.width / 8;
|
||||
uint32_t unpacked_mbr_size;
|
||||
if (mbr_type.vecsize == 3)
|
||||
unpacked_mbr_size = component_size * (mbr_type.vecsize + 1) * mbr_type.columns;
|
||||
else
|
||||
unpacked_mbr_size = component_size * mbr_type.vecsize * mbr_type.columns;
|
||||
if (is_array(mbr_type))
|
||||
{
|
||||
// If member is an array, and the array stride is larger than the type needs, don't pack it.
|
||||
|
|
@ -1640,11 +1659,49 @@ MSLStructMemberKey CompilerMSL::get_struct_member_key(uint32_t type_id, uint32_t
|
|||
return k;
|
||||
}
|
||||
|
||||
void CompilerMSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression)
|
||||
{
|
||||
if (!has_extended_decoration(lhs_expression, SPIRVCrossDecorationPacked) ||
|
||||
get_extended_decoration(lhs_expression, SPIRVCrossDecorationPackedType) == 0)
|
||||
{
|
||||
CompilerGLSL::emit_store_statement(lhs_expression, rhs_expression);
|
||||
}
|
||||
else
|
||||
{
|
||||
// Special handling when storing to a float[] or float2[] in std140 layout.
|
||||
|
||||
auto &type = get<SPIRType>(get_extended_decoration(lhs_expression, SPIRVCrossDecorationPackedType));
|
||||
string lhs = to_dereferenced_expression(lhs_expression);
|
||||
string rhs = to_pointer_expression(rhs_expression);
|
||||
|
||||
// Unpack the expression so we can store to it with a float or float2.
|
||||
// It's still an l-value, so it's fine. Most other unpacking of expressions turn them into r-values instead.
|
||||
if (is_scalar(type) && is_array(type))
|
||||
lhs = enclose_expression(lhs) + ".x";
|
||||
else if (is_vector(type) && type.vecsize == 2 && is_array(type))
|
||||
lhs = enclose_expression(lhs) + ".xy";
|
||||
|
||||
if (!optimize_read_modify_write(expression_type(rhs_expression), lhs, rhs))
|
||||
statement(lhs, " = ", rhs, ";");
|
||||
register_write(lhs_expression);
|
||||
}
|
||||
}
|
||||
|
||||
// Converts the format of the current expression from packed to unpacked,
|
||||
// by wrapping the expression in a constructor of the appropriate type.
|
||||
string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type)
|
||||
string CompilerMSL::unpack_expression_type(string expr_str, const SPIRType &type, uint32_t packed_type_id)
|
||||
{
|
||||
return join(type_to_glsl(type), "(", expr_str, ")");
|
||||
const SPIRType *packed_type = nullptr;
|
||||
if (packed_type_id)
|
||||
packed_type = &get<SPIRType>(packed_type_id);
|
||||
|
||||
// float[] and float2[] cases are really just padding, so directly swizzle from the backing float4 instead.
|
||||
if (packed_type && is_array(*packed_type) && is_scalar(*packed_type))
|
||||
return enclose_expression(expr_str) + ".x";
|
||||
else if (packed_type && is_array(*packed_type) && is_vector(*packed_type) && packed_type->vecsize == 2)
|
||||
return enclose_expression(expr_str) + ".xy";
|
||||
else
|
||||
return join(type_to_glsl(type), "(", expr_str, ")");
|
||||
}
|
||||
|
||||
// Emits the file header info
|
||||
|
|
@ -2109,8 +2166,9 @@ void CompilerMSL::emit_custom_functions()
|
|||
statement("");
|
||||
statement("// Wrapper function that swizzles texture gathers.");
|
||||
statement("template<typename T, typename Tex, typename... Ts>");
|
||||
statement("inline vec<T, 4> spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) "
|
||||
"METAL_CONST_ARG(c)");
|
||||
statement(
|
||||
"inline vec<T, 4> spvGatherSwizzle(sampler s, const thread Tex& t, Ts... params, component c, uint sw) "
|
||||
"METAL_CONST_ARG(c)");
|
||||
begin_scope();
|
||||
statement("if (sw)");
|
||||
begin_scope();
|
||||
|
|
@ -2149,7 +2207,8 @@ void CompilerMSL::emit_custom_functions()
|
|||
statement("");
|
||||
statement("// Wrapper function that swizzles depth texture gathers.");
|
||||
statement("template<typename T, typename Tex, typename... Ts>");
|
||||
statement("inline vec<T, 4> spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) ");
|
||||
statement(
|
||||
"inline vec<T, 4> spvGatherCompareSwizzle(sampler s, const thread Tex& t, Ts... params, uint sw) ");
|
||||
begin_scope();
|
||||
statement("if (sw)");
|
||||
begin_scope();
|
||||
|
|
@ -2330,7 +2389,7 @@ void CompilerMSL::emit_specialization_constants_and_structs()
|
|||
|
||||
declared_structs.insert(type_id);
|
||||
|
||||
if (has_decoration(type_id, DecorationCPacked))
|
||||
if (has_extended_decoration(type_id, SPIRVCrossDecorationPacked))
|
||||
align_struct(type);
|
||||
|
||||
// Make sure we declare the underlying struct type, and not the "decorated" type with pointers, etc.
|
||||
|
|
@ -2652,8 +2711,13 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
|
|||
test(bias, ImageOperandsBiasMask);
|
||||
test(lod, ImageOperandsLodMask);
|
||||
|
||||
auto &texel_type = expression_type(texel_id);
|
||||
auto store_type = texel_type;
|
||||
store_type.vecsize = 4;
|
||||
|
||||
statement(join(
|
||||
to_expression(img_id), ".write(", to_expression(texel_id), ", ",
|
||||
to_expression(img_id), ".write(", remap_swizzle(store_type, texel_type.vecsize, to_expression(texel_id)),
|
||||
", ",
|
||||
to_function_args(img_id, img_type, true, false, false, coord_id, 0, 0, 0, 0, lod, 0, 0, 0, 0, 0, &forward),
|
||||
");"));
|
||||
|
||||
|
|
@ -2813,7 +2877,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
|
|||
uint32_t mtx_id = ops[opcode == OpMatrixTimesVector ? 2 : 3];
|
||||
auto *e = maybe_get<SPIRExpression>(mtx_id);
|
||||
auto &t = expression_type(mtx_id);
|
||||
bool is_packed = has_decoration(mtx_id, DecorationCPacked);
|
||||
bool is_packed = has_extended_decoration(mtx_id, SPIRVCrossDecorationPacked);
|
||||
if (e && e->need_transpose && (t.columns == t.vecsize || is_packed))
|
||||
{
|
||||
e->need_transpose = false;
|
||||
|
|
@ -2822,11 +2886,11 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
|
|||
// are generally transposed, so unpacking using a constructor argument
|
||||
// will result in an error.
|
||||
// The simplest solution for now is to just avoid unpacking the matrix in this operation.
|
||||
unset_decoration(mtx_id, DecorationCPacked);
|
||||
unset_extended_decoration(mtx_id, SPIRVCrossDecorationPacked);
|
||||
|
||||
emit_binary_op(ops[0], ops[1], ops[3], ops[2], "*");
|
||||
if (is_packed)
|
||||
set_decoration(mtx_id, DecorationCPacked);
|
||||
set_extended_decoration(mtx_id, SPIRVCrossDecorationPacked);
|
||||
e->need_transpose = true;
|
||||
}
|
||||
else
|
||||
|
|
@ -3859,7 +3923,7 @@ bool CompilerMSL::is_non_native_row_major_matrix(uint32_t id)
|
|||
|
||||
// Generate a function that will swap matrix elements from row-major to column-major.
|
||||
// Packed row-matrix should just use transpose() function.
|
||||
if (!has_decoration(id, DecorationCPacked))
|
||||
if (!has_extended_decoration(id, SPIRVCrossDecorationPacked))
|
||||
{
|
||||
const auto type = expression_type(id);
|
||||
add_convert_row_major_matrix_function(type.columns, type.vecsize);
|
||||
|
|
@ -3881,7 +3945,7 @@ bool CompilerMSL::member_is_non_native_row_major_matrix(const SPIRType &type, ui
|
|||
|
||||
// Generate a function that will swap matrix elements from row-major to column-major.
|
||||
// Packed row-matrix should just use transpose() function.
|
||||
if (!has_member_decoration(type.self, index, DecorationCPacked))
|
||||
if (!has_extended_member_decoration(type.self, index, SPIRVCrossDecorationPacked))
|
||||
{
|
||||
const auto mbr_type = get<SPIRType>(type.member_types[index]);
|
||||
add_convert_row_major_matrix_function(mbr_type.columns, mbr_type.vecsize);
|
||||
|
|
@ -3964,13 +4028,16 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_
|
|||
|
||||
// If this member is packed, mark it as so.
|
||||
string pack_pfx = "";
|
||||
|
||||
const SPIRType *effective_membertype = &membertype;
|
||||
SPIRType override_type;
|
||||
|
||||
if (member_is_packed_type(type, index))
|
||||
{
|
||||
pack_pfx = "packed_";
|
||||
|
||||
// If we're packing a matrix, output an appropriate typedef
|
||||
if (membertype.vecsize > 1 && membertype.columns > 1)
|
||||
{
|
||||
pack_pfx = "packed_";
|
||||
string base_type = membertype.width == 16 ? "half" : "float";
|
||||
string td_line = "typedef ";
|
||||
td_line += base_type + to_string(membertype.vecsize) + "x" + to_string(membertype.columns);
|
||||
|
|
@ -3979,10 +4046,19 @@ void CompilerMSL::emit_struct_member(const SPIRType &type, uint32_t member_type_
|
|||
td_line += ";";
|
||||
add_typedef_line(td_line);
|
||||
}
|
||||
else if (is_array(membertype) && membertype.vecsize <= 2 && membertype.basetype != SPIRType::Struct)
|
||||
{
|
||||
// A "packed" float array, but we pad here instead to 4-vector.
|
||||
override_type = membertype;
|
||||
override_type.vecsize = 4;
|
||||
effective_membertype = &override_type;
|
||||
}
|
||||
else
|
||||
pack_pfx = "packed_";
|
||||
}
|
||||
|
||||
statement(pack_pfx, type_to_glsl(membertype), " ", qualifier, to_member_name(type, index),
|
||||
member_attribute_qualifier(type, index), type_to_array_glsl(membertype), ";");
|
||||
statement(pack_pfx, type_to_glsl(*effective_membertype), " ", qualifier, to_member_name(type, index),
|
||||
member_attribute_qualifier(type, index), type_to_array_glsl(*effective_membertype), ";");
|
||||
}
|
||||
|
||||
// Return a MSL qualifier for the specified function attribute member
|
||||
|
|
@ -5576,7 +5652,7 @@ size_t CompilerMSL::get_declared_struct_member_size(const SPIRType &struct_type,
|
|||
uint32_t columns = type.columns;
|
||||
|
||||
// An unpacked 3-element vector or matrix column is the same memory size as a 4-element.
|
||||
if (vecsize == 3 && !has_member_decoration(struct_type.self, index, DecorationCPacked))
|
||||
if (vecsize == 3 && !has_extended_member_decoration(struct_type.self, index, SPIRVCrossDecorationPacked))
|
||||
vecsize = 4;
|
||||
|
||||
return component_size * vecsize * columns;
|
||||
|
|
|
|||
|
|
@ -346,7 +346,7 @@ protected:
|
|||
uint32_t grad_y, uint32_t lod, uint32_t coffset, uint32_t offset, uint32_t bias,
|
||||
uint32_t comp, uint32_t sample, bool *p_forward) override;
|
||||
std::string to_initializer_expression(const SPIRVariable &var) override;
|
||||
std::string unpack_expression_type(std::string expr_str, const SPIRType &type) override;
|
||||
std::string unpack_expression_type(std::string expr_str, const SPIRType &type, uint32_t packed_type_id) override;
|
||||
std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override;
|
||||
bool skip_argument(uint32_t id) const override;
|
||||
std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override;
|
||||
|
|
@ -430,6 +430,7 @@ protected:
|
|||
|
||||
void bitcast_to_builtin_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override;
|
||||
void bitcast_from_builtin_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) override;
|
||||
void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression) override;
|
||||
|
||||
void analyze_sampled_image_usage();
|
||||
|
||||
|
|
|
|||
Loading…
Reference in New Issue