mirror of
https://github.com/KhronosGroup/SPIRV-Cross.git
synced 2024-09-19 22:30:03 +00:00
Merge pull request #2062 from cdavis5e/msl-spirv-assembly-fixes
MSL: SPIR-V assembly fixes
This commit is contained in:
Hans-Kristian Arntzen
GitHub
commit
451ed10e91
@ -1,11 +1,11 @@
|
||||
static float3x4 _362;
|
||||
static float4x3 _364;
|
||||
|
||||
RWByteAddressBuffer _17 : register(u0);
|
||||
uniform RaytracingAccelerationStructure rtas : register(t1);
|
||||
|
||||
static RayQuery<RAY_FLAG_NONE> rayQuery;
|
||||
|
||||
static float3x4 _362;
|
||||
static float4x3 _364;
|
||||
|
||||
void comp_main()
|
||||
{
|
||||
RayDesc _1ident = {0.0f.xxx, 0.0f, float3(1.0f, 0.0f, 0.0f), 9999.0f};
|
||||
|
||||
@ -1,3 +1,5 @@
|
||||
static bool _47;
|
||||
|
||||
static float2 value;
|
||||
static float4 FragColor;
|
||||
|
||||
@ -11,8 +13,6 @@ struct SPIRV_Cross_Output
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
static bool _47;
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
bool2 _25 = bool2(value.x == 0.0f, _47);
|
||||
|
||||
@ -0,0 +1,23 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _19
|
||||
{
|
||||
};
|
||||
struct _5
|
||||
{
|
||||
int _m0;
|
||||
_19 _m1;
|
||||
char _m2_pad[4];
|
||||
_19 _m2;
|
||||
char _m3_pad[4];
|
||||
int _m3;
|
||||
};
|
||||
|
||||
kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
|
||||
{
|
||||
_4 = _3;
|
||||
}
|
||||
|
||||
@ -0,0 +1,24 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _19
|
||||
{
|
||||
};
|
||||
struct _5
|
||||
{
|
||||
int _m0;
|
||||
char _m1_pad[12];
|
||||
_19 _m1;
|
||||
char _m2_pad[16];
|
||||
_19 _m2;
|
||||
char _m3_pad[16];
|
||||
int _m3;
|
||||
};
|
||||
|
||||
kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
|
||||
{
|
||||
_4 = _3;
|
||||
}
|
||||
|
||||
@ -54,6 +54,7 @@ struct _7
|
||||
int _m0[1];
|
||||
};
|
||||
|
||||
constant int3 _32 = {};
|
||||
constant int _3_tmp [[function_constant(0)]];
|
||||
constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0;
|
||||
constant int _4_tmp [[function_constant(1)]];
|
||||
@ -102,8 +103,6 @@ constant int _74 = (_73 * _72);
|
||||
constant spvUnsafeArray<int, 3> _33 = spvUnsafeArray<int, 3>({ 0, 0, 0 });
|
||||
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _34 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
|
||||
|
||||
constant int3 _32 = {};
|
||||
|
||||
kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74);
|
||||
|
||||
@ -0,0 +1,25 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _4
|
||||
{
|
||||
uint _m0[1];
|
||||
};
|
||||
|
||||
struct _20
|
||||
{
|
||||
uint _m0;
|
||||
uint _m1;
|
||||
};
|
||||
|
||||
kernel void main0(device _4& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]], device _4& _7 [[buffer(2)]], device _4& _8 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_20 _28;
|
||||
_28._m0 = uint(int(_5._m0[gl_GlobalInvocationID.x]) * int(_6._m0[gl_GlobalInvocationID.x]));
|
||||
_28._m1 = uint(mulhi(int(_5._m0[gl_GlobalInvocationID.x]), int(_6._m0[gl_GlobalInvocationID.x])));
|
||||
_7._m0[gl_GlobalInvocationID.x] = _28._m0;
|
||||
_8._m0[gl_GlobalInvocationID.x] = _28._m1;
|
||||
}
|
||||
|
||||
@ -0,0 +1,28 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _20
|
||||
{
|
||||
int _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
struct _5
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
struct _7
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
constant int _28 = {};
|
||||
|
||||
kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_6._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + (_20{ _28, 200 })._m1;
|
||||
}
|
||||
|
||||
@ -0,0 +1,31 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _21
|
||||
{
|
||||
int _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
struct _5
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
struct _7
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
constant int _29 = {};
|
||||
constant int _9_tmp [[function_constant(0)]];
|
||||
constant int _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 0;
|
||||
constant _21 _30 = _21{ _9, _29 };
|
||||
|
||||
kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_6._m0[gl_GlobalInvocationID.x] = (_8._m0[gl_GlobalInvocationID.x] + _30._m0) + (_21{ _29, 200 })._m1;
|
||||
}
|
||||
|
||||
@ -16,6 +16,8 @@ struct type_Globals
|
||||
uint2 ShadowTileListGroupSize;
|
||||
};
|
||||
|
||||
constant float3 _70 = {};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
const device type_StructuredBuffer_v4float* CulledObjectBoxBounds [[id(0)]];
|
||||
@ -24,8 +26,6 @@ struct spvDescriptorSetBuffer0
|
||||
device atomic_uint* RWShadowTileNumCulledObjects_atomic [[id(3)]];
|
||||
};
|
||||
|
||||
constant float3 _70 = {};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 out_var_SV_Target0 [[color(0)]];
|
||||
|
||||
@ -0,0 +1,26 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
struct _20
|
||||
{
|
||||
int _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
int _28;
|
||||
|
||||
layout(binding = 1, std430) buffer _5_6
|
||||
{
|
||||
int _m0[10];
|
||||
} _6;
|
||||
|
||||
layout(binding = 0, std430) buffer _7_8
|
||||
{
|
||||
int _m0[10];
|
||||
} _8;
|
||||
|
||||
void main()
|
||||
{
|
||||
_6._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + _20(_28, 200)._m1;
|
||||
}
|
||||
|
||||
@ -1,9 +1,9 @@
|
||||
#version 450
|
||||
|
||||
layout(location = 0) out vec4 _3;
|
||||
|
||||
float _49;
|
||||
|
||||
layout(location = 0) out vec4 _3;
|
||||
|
||||
void main()
|
||||
{
|
||||
_3 = vec4(_49);
|
||||
|
||||
@ -12,6 +12,8 @@
|
||||
#define SPIRV_CROSS_LOOP
|
||||
#endif
|
||||
|
||||
int _231;
|
||||
|
||||
layout(binding = 0, std140) uniform Foo
|
||||
{
|
||||
layout(row_major) mat4 lightVP[64];
|
||||
@ -22,8 +24,6 @@ layout(binding = 0, std140) uniform Foo
|
||||
layout(location = 0) in vec3 fragWorld;
|
||||
layout(location = 0) out int _entryPointOutput;
|
||||
|
||||
int _231;
|
||||
|
||||
mat4 spvWorkaroundRowMajor(mat4 wrap) { return wrap; }
|
||||
|
||||
void main()
|
||||
|
||||
@ -8,12 +8,12 @@ struct Foo
|
||||
float var2;
|
||||
};
|
||||
|
||||
Foo _22;
|
||||
|
||||
layout(binding = 0) uniform mediump sampler2D uSampler;
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
Foo _22;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = texture(uSampler, vec2(_22.var1, _22.var2));
|
||||
|
||||
@ -1,13 +1,13 @@
|
||||
#version 310 es
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
float _188;
|
||||
|
||||
layout(binding = 0, std430) buffer SSBO
|
||||
{
|
||||
float data;
|
||||
} _11;
|
||||
|
||||
float _188;
|
||||
|
||||
void main()
|
||||
{
|
||||
if (_11.data != 0.0)
|
||||
|
||||
@ -2,10 +2,10 @@
|
||||
precision mediump float;
|
||||
precision highp int;
|
||||
|
||||
layout(location = 0) out vec4 fragColor;
|
||||
|
||||
vec4 _80;
|
||||
|
||||
layout(location = 0) out vec4 fragColor;
|
||||
|
||||
void main()
|
||||
{
|
||||
mediump int _18 = int(_80.x);
|
||||
|
||||
@ -1,7 +1,7 @@
|
||||
RWByteAddressBuffer block : register(u0);
|
||||
|
||||
static float _15;
|
||||
|
||||
RWByteAddressBuffer block : register(u0);
|
||||
|
||||
void comp_main()
|
||||
{
|
||||
block.Store4(0, asuint(float4(0.100000001490116119384765625f, 0.20000000298023223876953125f, 0.300000011920928955078125f, 0.0f)));
|
||||
|
||||
@ -1,3 +1,5 @@
|
||||
static float4 _32;
|
||||
|
||||
static const float4 _34[2] = { 0.0f.xxxx, 0.0f.xxxx };
|
||||
|
||||
static float4 vInput;
|
||||
@ -13,8 +15,6 @@ struct SPIRV_Cross_Output
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
static float4 _32;
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
float4 _37 = vInput;
|
||||
|
||||
@ -1,8 +1,14 @@
|
||||
static int uninit_int = 0;
|
||||
static int4 uninit_vector = int4(0, 0, 0, 0);
|
||||
static float4x4 uninit_matrix = float4x4(0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx);
|
||||
|
||||
struct Foo
|
||||
{
|
||||
int a;
|
||||
};
|
||||
|
||||
static Foo uninit_foo = { 0 };
|
||||
|
||||
static float4 vColor;
|
||||
static float4 FragColor;
|
||||
|
||||
@ -16,11 +22,6 @@ struct SPIRV_Cross_Output
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
static int uninit_int = 0;
|
||||
static int4 uninit_vector = int4(0, 0, 0, 0);
|
||||
static float4x4 uninit_matrix = float4x4(0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx, 0.0f.xxxx);
|
||||
static Foo uninit_foo = { 0 };
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
int _39 = 0;
|
||||
|
||||
@ -1,3 +1,5 @@
|
||||
static float4 undef;
|
||||
|
||||
static float4 FragColor;
|
||||
static float4 vFloat;
|
||||
|
||||
@ -11,8 +13,6 @@ struct SPIRV_Cross_Output
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
static float4 undef;
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
FragColor = float4(undef.x, vFloat.y, 0.0f, vFloat.w) + float4(vFloat.z, vFloat.y, 0.0f, vFloat.w);
|
||||
|
||||
@ -1,3 +1,5 @@
|
||||
static float4 _21;
|
||||
|
||||
static int counter;
|
||||
static float4 FragColor;
|
||||
|
||||
@ -11,8 +13,6 @@ struct SPIRV_Cross_Output
|
||||
float4 FragColor : SV_Target0;
|
||||
};
|
||||
|
||||
static float4 _21;
|
||||
|
||||
void frag_main()
|
||||
{
|
||||
float4 _24;
|
||||
|
||||
@ -44,10 +44,10 @@ struct spvUnsafeArray
|
||||
}
|
||||
};
|
||||
|
||||
constant spvUnsafeArray<float4, 2> _34 = spvUnsafeArray<float4, 2>({ float4(0.0), float4(0.0) });
|
||||
|
||||
constant float4 _32 = {};
|
||||
|
||||
constant spvUnsafeArray<float4, 2> _34 = spvUnsafeArray<float4, 2>({ float4(0.0), float4(0.0) });
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 FragColor [[color(0)]];
|
||||
|
||||
@ -3,14 +3,15 @@
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant int uninit_int = {};
|
||||
constant int4 uninit_vector = {};
|
||||
constant float4x4 uninit_matrix = {};
|
||||
|
||||
struct Foo
|
||||
{
|
||||
int a;
|
||||
};
|
||||
|
||||
constant int uninit_int = {};
|
||||
constant int4 uninit_vector = {};
|
||||
constant float4x4 uninit_matrix = {};
|
||||
constant Foo uninit_foo = {};
|
||||
|
||||
struct main0_out
|
||||
|
||||
@ -0,0 +1,23 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _19
|
||||
{
|
||||
};
|
||||
struct _5
|
||||
{
|
||||
int _m0;
|
||||
_19 _m1;
|
||||
char _m2_pad[4];
|
||||
_19 _m2;
|
||||
char _m3_pad[4];
|
||||
int _m3;
|
||||
};
|
||||
|
||||
kernel void main0(device _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
|
||||
{
|
||||
_4 = _3;
|
||||
}
|
||||
|
||||
@ -0,0 +1,24 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _19
|
||||
{
|
||||
};
|
||||
struct _5
|
||||
{
|
||||
int _m0;
|
||||
char _m1_pad[12];
|
||||
_19 _m1;
|
||||
char _m2_pad[16];
|
||||
_19 _m2;
|
||||
char _m3_pad[16];
|
||||
int _m3;
|
||||
};
|
||||
|
||||
kernel void main0(constant _5& _3 [[buffer(0)]], device _5& _4 [[buffer(1)]])
|
||||
{
|
||||
_4 = _3;
|
||||
}
|
||||
|
||||
@ -54,6 +54,7 @@ struct _7
|
||||
int _m0[1];
|
||||
};
|
||||
|
||||
constant int3 _32 = {};
|
||||
constant int _3_tmp [[function_constant(0)]];
|
||||
constant int _3 = is_function_constant_defined(_3_tmp) ? _3_tmp : 0;
|
||||
constant int _4_tmp [[function_constant(1)]];
|
||||
@ -103,8 +104,6 @@ constant int _74 = (_73 * _72);
|
||||
constant spvUnsafeArray<int, 3> _33 = spvUnsafeArray<int, 3>({ 0, 0, 0 });
|
||||
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _34 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
|
||||
|
||||
constant int3 _32 = {};
|
||||
|
||||
kernel void main0(device _7& _8 [[buffer(0)]], device _7& _9 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_9._m0[gl_GlobalInvocationID.x] = _8._m0[gl_GlobalInvocationID.x] + ((((1 - _59) * _60) * (_61 - 1)) * _74);
|
||||
|
||||
25
reference/shaders-msl/asm/comp/uint_smulextended.asm.comp
Normal file
25
reference/shaders-msl/asm/comp/uint_smulextended.asm.comp
Normal file
@ -0,0 +1,25 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _4
|
||||
{
|
||||
uint _m0[1];
|
||||
};
|
||||
|
||||
struct _20
|
||||
{
|
||||
uint _m0;
|
||||
uint _m1;
|
||||
};
|
||||
|
||||
kernel void main0(device _4& _5 [[buffer(0)]], device _4& _6 [[buffer(1)]], device _4& _7 [[buffer(2)]], device _4& _8 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
_20 _28;
|
||||
_28._m0 = uint(int(_5._m0[gl_GlobalInvocationID.x]) * int(_6._m0[gl_GlobalInvocationID.x]));
|
||||
_28._m1 = uint(mulhi(int(_5._m0[gl_GlobalInvocationID.x]), int(_6._m0[gl_GlobalInvocationID.x])));
|
||||
_7._m0[gl_GlobalInvocationID.x] = _28._m0;
|
||||
_8._m0[gl_GlobalInvocationID.x] = _28._m1;
|
||||
}
|
||||
|
||||
@ -0,0 +1,38 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _20
|
||||
{
|
||||
int _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
struct _5
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
struct _7
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
constant int _28 = {};
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
int _39(thread const int& _41, thread const _20& _42)
|
||||
{
|
||||
return _41 + _42._m1;
|
||||
}
|
||||
|
||||
kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
int _32 = _8._m0[gl_GlobalInvocationID.x];
|
||||
_20 _33 = _20{ _28, 200 };
|
||||
_6._m0[gl_GlobalInvocationID.x] = _39(_32, _33);
|
||||
}
|
||||
|
||||
@ -0,0 +1,42 @@
|
||||
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
struct _21
|
||||
{
|
||||
int _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
struct _5
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
struct _7
|
||||
{
|
||||
int _m0[10];
|
||||
};
|
||||
|
||||
constant int _29 = {};
|
||||
constant int _9_tmp [[function_constant(0)]];
|
||||
constant int _9 = is_function_constant_defined(_9_tmp) ? _9_tmp : 0;
|
||||
constant _21 _30 = _21{ _9, _29 };
|
||||
|
||||
static inline __attribute__((always_inline))
|
||||
int _42(thread const int& _44, thread const _21& _45, thread const _21& _46)
|
||||
{
|
||||
return (_44 + _45._m0) + _46._m1;
|
||||
}
|
||||
|
||||
kernel void main0(device _5& _6 [[buffer(0)]], device _7& _8 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
int _34 = _8._m0[gl_GlobalInvocationID.x];
|
||||
_21 _35 = _30;
|
||||
_21 _36 = _21{ _29, 200 };
|
||||
_6._m0[gl_GlobalInvocationID.x] = _42(_34, _35, _36);
|
||||
}
|
||||
|
||||
@ -1,13 +1,13 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
float _15;
|
||||
|
||||
layout(binding = 0, std430) buffer Block
|
||||
{
|
||||
vec4 f;
|
||||
} block;
|
||||
|
||||
float _15;
|
||||
|
||||
void main()
|
||||
{
|
||||
block.f = vec4(0.100000001490116119384765625, 0.20000000298023223876953125, 0.300000011920928955078125, 0.0);
|
||||
|
||||
@ -2,11 +2,11 @@
|
||||
#extension GL_EXT_ray_query : require
|
||||
layout(local_size_x = 64, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
layout(set = 0, binding = 0) uniform accelerationStructureEXT RTAS;
|
||||
|
||||
float _16;
|
||||
vec3 _17;
|
||||
|
||||
layout(set = 0, binding = 0) uniform accelerationStructureEXT RTAS;
|
||||
|
||||
void main()
|
||||
{
|
||||
rayQueryEXT _19;
|
||||
|
||||
@ -2,11 +2,11 @@
|
||||
precision mediump float;
|
||||
precision highp int;
|
||||
|
||||
vec4 _32;
|
||||
|
||||
layout(location = 0) in vec4 vInput;
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
vec4 _32;
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 _37 = vInput;
|
||||
|
||||
@ -1,5 +1,7 @@
|
||||
#version 450
|
||||
|
||||
vec4 _32;
|
||||
|
||||
layout(binding = 0, std140) uniform type_gCBuffarrayIndex
|
||||
{
|
||||
uint gArrayIndex;
|
||||
@ -12,8 +14,6 @@ uniform sampler2D SPIRV_Cross_Combinedg_textureArray3SPIRV_Cross_DummySampler;
|
||||
|
||||
layout(location = 0) out vec4 out_var_SV_TARGET;
|
||||
|
||||
vec4 _32;
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 _80;
|
||||
|
||||
@ -61,6 +61,15 @@ struct Params
|
||||
vec4 LqmatFarTilingFactor;
|
||||
};
|
||||
|
||||
VertexOutput _121;
|
||||
SurfaceInput _122;
|
||||
vec2 _123;
|
||||
vec4 _124;
|
||||
Surface _125;
|
||||
vec4 _192;
|
||||
vec4 _219;
|
||||
vec4 _297;
|
||||
|
||||
layout(binding = 0, std140) uniform CB0
|
||||
{
|
||||
Globals CB0;
|
||||
@ -86,15 +95,6 @@ layout(location = 7) in vec4 IN_PosLightSpace_Reflectance;
|
||||
layout(location = 8) in float IN_studIndex;
|
||||
layout(location = 0) out vec4 _entryPointOutput;
|
||||
|
||||
VertexOutput _121;
|
||||
SurfaceInput _122;
|
||||
vec2 _123;
|
||||
vec4 _124;
|
||||
Surface _125;
|
||||
vec4 _192;
|
||||
vec4 _219;
|
||||
vec4 _297;
|
||||
|
||||
void main()
|
||||
{
|
||||
VertexOutput _128;
|
||||
|
||||
@ -1,18 +1,19 @@
|
||||
#version 450
|
||||
|
||||
int uninit_int = 0;
|
||||
ivec4 uninit_vector = ivec4(0);
|
||||
mat4 uninit_matrix = mat4(vec4(0.0), vec4(0.0), vec4(0.0), vec4(0.0));
|
||||
|
||||
struct Foo
|
||||
{
|
||||
int a;
|
||||
};
|
||||
|
||||
Foo uninit_foo = Foo(0);
|
||||
|
||||
layout(location = 0) in vec4 vColor;
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
int uninit_int = 0;
|
||||
ivec4 uninit_vector = ivec4(0);
|
||||
mat4 uninit_matrix = mat4(vec4(0.0), vec4(0.0), vec4(0.0), vec4(0.0));
|
||||
Foo uninit_foo = Foo(0);
|
||||
|
||||
void main()
|
||||
{
|
||||
int _39 = 0;
|
||||
|
||||
@ -12,12 +12,12 @@ struct _5
|
||||
int _m1;
|
||||
};
|
||||
|
||||
layout(location = 0) flat in int _2;
|
||||
layout(location = 0) out int _3;
|
||||
|
||||
_4 _16;
|
||||
int _21;
|
||||
|
||||
layout(location = 0) flat in int _2;
|
||||
layout(location = 0) out int _3;
|
||||
|
||||
void main()
|
||||
{
|
||||
bool _25 = false;
|
||||
|
||||
@ -2,10 +2,10 @@
|
||||
precision mediump float;
|
||||
precision highp int;
|
||||
|
||||
layout(location = 0) out highp vec4 _GLF_color;
|
||||
|
||||
vec2 _19;
|
||||
|
||||
layout(location = 0) out highp vec4 _GLF_color;
|
||||
|
||||
void main()
|
||||
{
|
||||
highp vec2 _30;
|
||||
|
||||
@ -1,10 +1,10 @@
|
||||
#version 450
|
||||
|
||||
vec4 undef;
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
layout(location = 0) in vec4 vFloat;
|
||||
|
||||
vec4 undef;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = vec4(undef.x, vFloat.y, 0.0, vFloat.w) + vec4(vFloat.z, vFloat.y, 0.0, vFloat.w);
|
||||
|
||||
@ -93,6 +93,8 @@ struct type_Globals
|
||||
float ExpandGamut;
|
||||
};
|
||||
|
||||
constant float3 _391 = {};
|
||||
|
||||
constant spvUnsafeArray<float, 6> _475 = spvUnsafeArray<float, 6>({ -4.0, -4.0, -3.1573765277862548828125, -0.485249996185302734375, 1.84773242473602294921875, 1.84773242473602294921875 });
|
||||
constant spvUnsafeArray<float, 6> _476 = spvUnsafeArray<float, 6>({ -0.718548238277435302734375, 2.0810306072235107421875, 3.66812419891357421875, 4.0, 4.0, 4.0 });
|
||||
constant spvUnsafeArray<float, 10> _479 = spvUnsafeArray<float, 10>({ -4.97062206268310546875, -3.0293781757354736328125, -2.1261999607086181640625, -1.5104999542236328125, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 });
|
||||
@ -100,8 +102,6 @@ constant spvUnsafeArray<float, 10> _480 = spvUnsafeArray<float, 10>({ 0.80891323
|
||||
constant spvUnsafeArray<float, 10> _482 = spvUnsafeArray<float, 10>({ -2.3010299205780029296875, -2.3010299205780029296875, -1.9312000274658203125, -1.5204999446868896484375, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 });
|
||||
constant spvUnsafeArray<float, 10> _483 = spvUnsafeArray<float, 10>({ 0.801995217800140380859375, 1.19800484180450439453125, 1.5943000316619873046875, 1.99730002880096435546875, 2.3782999515533447265625, 2.7683999538421630859375, 3.0515000820159912109375, 3.2746293544769287109375, 3.32743072509765625, 3.32743072509765625 });
|
||||
|
||||
constant float3 _391 = {};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 out_var_SV_Target0 [[color(0)]];
|
||||
|
||||
@ -94,6 +94,9 @@ struct type_Globals
|
||||
float ExpandGamut;
|
||||
};
|
||||
|
||||
constant float3 _523 = {};
|
||||
constant float3 _525 = {};
|
||||
|
||||
constant spvUnsafeArray<float, 6> _499 = spvUnsafeArray<float, 6>({ -4.0, -4.0, -3.1573765277862548828125, -0.485249996185302734375, 1.84773242473602294921875, 1.84773242473602294921875 });
|
||||
constant spvUnsafeArray<float, 6> _500 = spvUnsafeArray<float, 6>({ -0.718548238277435302734375, 2.0810306072235107421875, 3.66812419891357421875, 4.0, 4.0, 4.0 });
|
||||
constant spvUnsafeArray<float, 10> _503 = spvUnsafeArray<float, 10>({ -4.97062206268310546875, -3.0293781757354736328125, -2.1261999607086181640625, -1.5104999542236328125, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 });
|
||||
@ -101,9 +104,6 @@ constant spvUnsafeArray<float, 10> _504 = spvUnsafeArray<float, 10>({ 0.80891323
|
||||
constant spvUnsafeArray<float, 10> _506 = spvUnsafeArray<float, 10>({ -2.3010299205780029296875, -2.3010299205780029296875, -1.9312000274658203125, -1.5204999446868896484375, -1.0578000545501708984375, -0.4668000042438507080078125, 0.11937999725341796875, 0.7088134288787841796875, 1.2911865711212158203125, 1.2911865711212158203125 });
|
||||
constant spvUnsafeArray<float, 10> _507 = spvUnsafeArray<float, 10>({ 0.801995217800140380859375, 1.19800484180450439453125, 1.5943000316619873046875, 1.99730002880096435546875, 2.3782999515533447265625, 2.7683999538421630859375, 3.0515000820159912109375, 3.2746293544769287109375, 3.32743072509765625, 3.32743072509765625 });
|
||||
|
||||
constant float3 _523 = {};
|
||||
constant float3 _525 = {};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 out_var_SV_Target0 [[color(0)]];
|
||||
|
||||
@ -56,6 +56,8 @@ struct type_Globals
|
||||
uint2 ShadowTileListGroupSize;
|
||||
};
|
||||
|
||||
constant float3 _70 = {};
|
||||
|
||||
struct spvDescriptorSetBuffer0
|
||||
{
|
||||
const device type_StructuredBuffer_v4float* CulledObjectBoxBounds [[id(0)]];
|
||||
@ -64,8 +66,6 @@ struct spvDescriptorSetBuffer0
|
||||
device atomic_uint* RWShadowTileNumCulledObjects_atomic [[id(3)]];
|
||||
};
|
||||
|
||||
constant float3 _70 = {};
|
||||
|
||||
struct main0_out
|
||||
{
|
||||
float4 out_var_SV_Target0 [[color(0)]];
|
||||
|
||||
@ -0,0 +1,33 @@
|
||||
#version 450
|
||||
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
struct _20
|
||||
{
|
||||
int _m0;
|
||||
int _m1;
|
||||
};
|
||||
|
||||
int _28;
|
||||
|
||||
layout(binding = 1, std430) buffer _5_6
|
||||
{
|
||||
int _m0[10];
|
||||
} _6;
|
||||
|
||||
layout(binding = 0, std430) buffer _7_8
|
||||
{
|
||||
int _m0[10];
|
||||
} _8;
|
||||
|
||||
int _39(int _41, _20 _42)
|
||||
{
|
||||
return _41 + _42._m1;
|
||||
}
|
||||
|
||||
void main()
|
||||
{
|
||||
int _32 = _8._m0[gl_GlobalInvocationID.x];
|
||||
_20 _33 = _20(_28, 200);
|
||||
_6._m0[gl_GlobalInvocationID.x] = _39(_32, _33);
|
||||
}
|
||||
|
||||
@ -8,12 +8,12 @@ struct Foo
|
||||
float var2;
|
||||
};
|
||||
|
||||
Foo _22;
|
||||
|
||||
layout(binding = 0) uniform mediump sampler2D uSampler;
|
||||
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
Foo _22;
|
||||
|
||||
void main()
|
||||
{
|
||||
FragColor = texture(uSampler, vec2(_22.var1, _22.var2));
|
||||
|
||||
@ -1,10 +1,10 @@
|
||||
#version 450
|
||||
|
||||
layout(location = 0) out vec4 _entryPointOutput;
|
||||
|
||||
vec4 _38;
|
||||
vec4 _47;
|
||||
|
||||
layout(location = 0) out vec4 _entryPointOutput;
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 _27;
|
||||
|
||||
@ -1,10 +1,10 @@
|
||||
#version 450
|
||||
|
||||
vec4 _21;
|
||||
|
||||
layout(location = 0) flat in int counter;
|
||||
layout(location = 0) out vec4 FragColor;
|
||||
|
||||
vec4 _21;
|
||||
|
||||
void main()
|
||||
{
|
||||
vec4 _24;
|
||||
|
||||
@ -17,6 +17,8 @@ struct _28
|
||||
vec4 _m0;
|
||||
};
|
||||
|
||||
_28 _74;
|
||||
|
||||
layout(binding = 0, std140) uniform _6_7
|
||||
{
|
||||
vec4 _m0;
|
||||
@ -102,8 +104,6 @@ uniform sampler2D SPIRV_Cross_Combined_2;
|
||||
|
||||
layout(location = 0) out vec4 _5;
|
||||
|
||||
_28 _74;
|
||||
|
||||
void main()
|
||||
{
|
||||
_28 _77;
|
||||
|
||||
43
shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp
Normal file
43
shaders-msl/asm/comp/copy-object-ssbo-to-ssbo.asm.comp
Normal file
@ -0,0 +1,43 @@
|
||||
OpCapability Shader
|
||||
OpExtension "SPV_KHR_storage_buffer_storage_class"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %var_id
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpDecorate %var_id BuiltIn GlobalInvocationId
|
||||
OpDecorate %var_input Binding 0
|
||||
OpDecorate %var_input DescriptorSet 0
|
||||
OpDecorate %var_outdata Binding 1
|
||||
OpDecorate %var_outdata DescriptorSet 0
|
||||
OpMemberDecorate %type_container_struct 0 Offset 0
|
||||
OpMemberDecorate %type_container_struct 1 Offset 4
|
||||
OpMemberDecorate %type_container_struct 2 Offset 8
|
||||
OpMemberDecorate %type_container_struct 3 Offset 12
|
||||
OpDecorate %type_container_struct Block
|
||||
%bool = OpTypeBool
|
||||
%void = OpTypeVoid
|
||||
%voidf = OpTypeFunction %void
|
||||
%u32 = OpTypeInt 32 0
|
||||
%i32 = OpTypeInt 32 1
|
||||
%f32 = OpTypeFloat 32
|
||||
%uvec3 = OpTypeVector %u32 3
|
||||
%fvec3 = OpTypeVector %f32 3
|
||||
%uvec3ptr = OpTypePointer Input %uvec3
|
||||
%i32ptr = OpTypePointer Uniform %i32
|
||||
%f32ptr = OpTypePointer Uniform %f32
|
||||
%i32arr = OpTypeRuntimeArray %i32
|
||||
%f32arr = OpTypeRuntimeArray %f32
|
||||
%type_empty_struct = OpTypeStruct
|
||||
%type_container_struct = OpTypeStruct %i32 %type_empty_struct %type_empty_struct %i32
|
||||
%type_container_struct_ubo_ptr = OpTypePointer Uniform %type_container_struct
|
||||
%type_container_struct_ssbo_ptr = OpTypePointer StorageBuffer %type_container_struct
|
||||
%var_id = OpVariable %uvec3ptr Input
|
||||
%var_input = OpVariable %type_container_struct_ssbo_ptr StorageBuffer
|
||||
%var_outdata = OpVariable %type_container_struct_ssbo_ptr StorageBuffer
|
||||
|
||||
%main = OpFunction %void None %voidf
|
||||
%label = OpLabel
|
||||
%input_copy = OpCopyObject %type_container_struct_ssbo_ptr %var_input
|
||||
%result = OpLoad %type_container_struct %input_copy
|
||||
OpStore %var_outdata %result
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
43
shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp
Normal file
43
shaders-msl/asm/comp/copy-object-ubo-to-ssbo.asm.comp
Normal file
@ -0,0 +1,43 @@
|
||||
OpCapability Shader
|
||||
OpExtension "SPV_KHR_storage_buffer_storage_class"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %var_id
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpDecorate %var_id BuiltIn GlobalInvocationId
|
||||
OpDecorate %var_input Binding 0
|
||||
OpDecorate %var_input DescriptorSet 0
|
||||
OpDecorate %var_outdata Binding 1
|
||||
OpDecorate %var_outdata DescriptorSet 0
|
||||
OpMemberDecorate %type_container_struct 0 Offset 0
|
||||
OpMemberDecorate %type_container_struct 1 Offset 16
|
||||
OpMemberDecorate %type_container_struct 2 Offset 32
|
||||
OpMemberDecorate %type_container_struct 3 Offset 48
|
||||
OpDecorate %type_container_struct Block
|
||||
%bool = OpTypeBool
|
||||
%void = OpTypeVoid
|
||||
%voidf = OpTypeFunction %void
|
||||
%u32 = OpTypeInt 32 0
|
||||
%i32 = OpTypeInt 32 1
|
||||
%f32 = OpTypeFloat 32
|
||||
%uvec3 = OpTypeVector %u32 3
|
||||
%fvec3 = OpTypeVector %f32 3
|
||||
%uvec3ptr = OpTypePointer Input %uvec3
|
||||
%i32ptr = OpTypePointer Uniform %i32
|
||||
%f32ptr = OpTypePointer Uniform %f32
|
||||
%i32arr = OpTypeRuntimeArray %i32
|
||||
%f32arr = OpTypeRuntimeArray %f32
|
||||
%type_empty_struct = OpTypeStruct
|
||||
%type_container_struct = OpTypeStruct %i32 %type_empty_struct %type_empty_struct %i32
|
||||
%type_container_struct_ubo_ptr = OpTypePointer Uniform %type_container_struct
|
||||
%type_container_struct_ssbo_ptr = OpTypePointer StorageBuffer %type_container_struct
|
||||
%var_id = OpVariable %uvec3ptr Input
|
||||
%var_input = OpVariable %type_container_struct_ubo_ptr Uniform
|
||||
%var_outdata = OpVariable %type_container_struct_ssbo_ptr StorageBuffer
|
||||
|
||||
%main = OpFunction %void None %voidf
|
||||
%label = OpLabel
|
||||
%input_copy = OpCopyObject %type_container_struct_ubo_ptr %var_input
|
||||
%result = OpLoad %type_container_struct %input_copy
|
||||
OpStore %var_outdata %result
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
61
shaders-msl/asm/comp/uint_smulextended.asm.comp
Normal file
61
shaders-msl/asm/comp/uint_smulextended.asm.comp
Normal file
@ -0,0 +1,61 @@
|
||||
OpCapability Shader
|
||||
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationId
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
|
||||
OpDecorate %gl_GlobalInvocationId BuiltIn GlobalInvocationId
|
||||
OpDecorate %ra_uint ArrayStride 4
|
||||
OpDecorate %struct_uint4 BufferBlock
|
||||
OpMemberDecorate %struct_uint4 0 Offset 0
|
||||
OpDecorate %input0 DescriptorSet 0
|
||||
OpDecorate %input0 Binding 0
|
||||
OpDecorate %input1 DescriptorSet 0
|
||||
OpDecorate %input1 Binding 1
|
||||
OpDecorate %output0 DescriptorSet 0
|
||||
OpDecorate %output0 Binding 2
|
||||
OpDecorate %output1 DescriptorSet 0
|
||||
OpDecorate %output1 Binding 3
|
||||
|
||||
%uint = OpTypeInt 32 0
|
||||
%ptr_uint = OpTypePointer Uniform %uint
|
||||
%ptr_input_uint = OpTypePointer Input %uint
|
||||
%uint3 = OpTypeVector %uint 3
|
||||
%ptr_input_uint3 = OpTypePointer Input %uint3
|
||||
%void = OpTypeVoid
|
||||
%voidFn = OpTypeFunction %void
|
||||
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%ra_uint = OpTypeRuntimeArray %uint
|
||||
%uint4 = OpTypeVector %uint 4
|
||||
%struct_uint4 = OpTypeStruct %ra_uint
|
||||
%ptr_struct_uint4 = OpTypePointer Uniform %struct_uint4
|
||||
%resulttype = OpTypeStruct %uint %uint
|
||||
%gl_GlobalInvocationId = OpVariable %ptr_input_uint3 Input
|
||||
%input0 = OpVariable %ptr_struct_uint4 Uniform
|
||||
%input1 = OpVariable %ptr_struct_uint4 Uniform
|
||||
|
||||
%output0 = OpVariable %ptr_struct_uint4 Uniform
|
||||
%output1 = OpVariable %ptr_struct_uint4 Uniform
|
||||
|
||||
%main = OpFunction %void None %voidFn
|
||||
%mainStart = OpLabel
|
||||
%index_ptr = OpAccessChain %ptr_input_uint %gl_GlobalInvocationId %uint_0
|
||||
%index = OpLoad %uint %index_ptr
|
||||
%in_ptr0 = OpAccessChain %ptr_uint %input0 %uint_0 %index
|
||||
%invalue0 = OpLoad %uint %in_ptr0
|
||||
%in_ptr1 = OpAccessChain %ptr_uint %input1 %uint_0 %index
|
||||
%invalue1 = OpLoad %uint %in_ptr1
|
||||
|
||||
%outvalue = OpSMulExtended %resulttype %invalue0 %invalue1
|
||||
%outvalue0 = OpCompositeExtract %uint %outvalue 0
|
||||
%out_ptr0 = OpAccessChain %ptr_uint %output0 %uint_0 %index
|
||||
OpStore %out_ptr0 %outvalue0
|
||||
%outvalue1 = OpCompositeExtract %uint %outvalue 1
|
||||
%out_ptr1 = OpAccessChain %ptr_uint %output1 %uint_0 %index
|
||||
OpStore %out_ptr1 %outvalue1
|
||||
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
102
shaders-msl/asm/comp/undefined-constant-composite.asm.comp
Normal file
102
shaders-msl/asm/comp/undefined-constant-composite.asm.comp
Normal file
@ -0,0 +1,102 @@
|
||||
;
|
||||
; The shader below is based on the following GLSL shader:
|
||||
;
|
||||
; #version 450
|
||||
;
|
||||
; struct Pair {
|
||||
; int first;
|
||||
; int second;
|
||||
; };
|
||||
;
|
||||
; const Pair constant_pair = { 100, 200 };
|
||||
;
|
||||
; layout(set=0, binding=0, std430) buffer InputBlock {
|
||||
; int array[10];
|
||||
; } inputValues;
|
||||
;
|
||||
; layout(set=0, binding=1, std430) buffer OutputBlock {
|
||||
; int array[10];
|
||||
; } outputValues;
|
||||
;
|
||||
; int add_second (int value, Pair pair) {
|
||||
; return value + pair.second;
|
||||
; }
|
||||
;
|
||||
; void main() {
|
||||
; uint idx = gl_GlobalInvocationID.x;
|
||||
; outputValues.array[idx] = add_second(inputValues.array[idx], constant_pair);
|
||||
; }
|
||||
;
|
||||
; However, the first element of constant_pair has been modified to be undefined.
|
||||
;
|
||||
OpCapability Shader
|
||||
%std450 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
|
||||
OpDecorate %_arr_int_uint_10 ArrayStride 4
|
||||
OpMemberDecorate %OutputBlock 0 Offset 0
|
||||
OpDecorate %OutputBlock BufferBlock
|
||||
OpDecorate %outputValues DescriptorSet 0
|
||||
OpDecorate %outputValues Binding 1
|
||||
OpMemberDecorate %InputBlock 0 Offset 0
|
||||
OpDecorate %InputBlock BufferBlock
|
||||
OpDecorate %inputValues DescriptorSet 0
|
||||
OpDecorate %inputValues Binding 0
|
||||
%void = OpTypeVoid
|
||||
%void_func = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%int_0 = OpConstant %int 0
|
||||
%int_1 = OpConstant %int 1
|
||||
%int_200 = OpConstant %int 200
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_10 = OpConstant %uint 10
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%Pair = OpTypeStruct %int %int
|
||||
%_ptr_Function_Pair = OpTypePointer Function %Pair
|
||||
%add_second_func_type = OpTypeFunction %int %_ptr_Function_int %_ptr_Function_Pair
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%_arr_int_uint_10 = OpTypeArray %int %uint_10
|
||||
%OutputBlock = OpTypeStruct %_arr_int_uint_10
|
||||
%_ptr_Uniform_OutputBlock = OpTypePointer Uniform %OutputBlock
|
||||
%outputValues = OpVariable %_ptr_Uniform_OutputBlock Uniform
|
||||
%InputBlock = OpTypeStruct %_arr_int_uint_10
|
||||
%_ptr_Uniform_InputBlock = OpTypePointer Uniform %InputBlock
|
||||
%inputValues = OpVariable %_ptr_Uniform_InputBlock Uniform
|
||||
; Replaced %int_100 with an undefined int.
|
||||
%undef_int = OpUndef %int
|
||||
; Composed a constant Pair with the undefined int in the first member.
|
||||
%const_Pair = OpConstantComposite %Pair %undef_int %int_200
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||
%main = OpFunction %void None %void_func
|
||||
%main_label = OpLabel
|
||||
%param_1 = OpVariable %_ptr_Function_int Function
|
||||
%param_2 = OpVariable %_ptr_Function_Pair Function
|
||||
%gidx_ptr = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||
%gidx = OpLoad %uint %gidx_ptr
|
||||
%input_value_ptr = OpAccessChain %_ptr_Uniform_int %inputValues %int_0 %gidx
|
||||
%input_value = OpLoad %int %input_value_ptr
|
||||
OpStore %param_1 %input_value
|
||||
OpStore %param_2 %const_Pair
|
||||
%retval = OpFunctionCall %int %add_second %param_1 %param_2
|
||||
%output_value_ptr = OpAccessChain %_ptr_Uniform_int %outputValues %int_0 %gidx
|
||||
OpStore %output_value_ptr %retval
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%add_second = OpFunction %int None %add_second_func_type
|
||||
%value_ptr = OpFunctionParameter %_ptr_Function_int
|
||||
%pair = OpFunctionParameter %_ptr_Function_Pair
|
||||
%add_second_label = OpLabel
|
||||
%value = OpLoad %int %value_ptr
|
||||
; Access the second struct member, which is defined.
|
||||
%pair_second_ptr = OpAccessChain %_ptr_Function_int %pair %int_1
|
||||
%pair_second = OpLoad %int %pair_second_ptr
|
||||
%add_result = OpIAdd %int %value %pair_second
|
||||
OpReturnValue %add_result
|
||||
OpFunctionEnd
|
||||
122
shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp
Normal file
122
shaders-msl/asm/comp/undefined-spec-constant-composite.asm.comp
Normal file
@ -0,0 +1,122 @@
|
||||
;
|
||||
; The shader below is based on the following GLSL shader:
|
||||
;
|
||||
; #version 450
|
||||
;
|
||||
; struct Pair {
|
||||
; int first;
|
||||
; int second;
|
||||
; };
|
||||
;
|
||||
; const Pair constant_pair = { 100, 200 };
|
||||
;
|
||||
; layout (constant_id=0) const int constantFirst = 0;
|
||||
;
|
||||
; Pair spec_constant_pair = { constantFirst, 200 };
|
||||
;
|
||||
; layout(set=0, binding=0, std430) buffer InputBlock {
|
||||
; int array[10];
|
||||
; } inputValues;
|
||||
;
|
||||
; layout(set=0, binding=1, std430) buffer OutputBlock {
|
||||
; int array[10];
|
||||
; } outputValues;
|
||||
;
|
||||
; int add_first_and_second (int value, Pair p1, Pair p2) {
|
||||
; return value + p1.first + p2.second;
|
||||
; }
|
||||
;
|
||||
; void main() {
|
||||
; uint idx = gl_GlobalInvocationID.x;
|
||||
; outputValues.array[idx] = add_first_and_second(inputValues.array[idx], spec_constant_pair, constant_pair);
|
||||
; }
|
||||
;
|
||||
; However, both the constant_pair and the spec_constant_pair have one of their members replaced by undefined values.
|
||||
;
|
||||
OpCapability Shader
|
||||
%std450 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
|
||||
OpDecorate %_arr_int_uint_10 ArrayStride 4
|
||||
OpMemberDecorate %OutputBlock 0 Offset 0
|
||||
OpDecorate %OutputBlock BufferBlock
|
||||
OpDecorate %outputValues DescriptorSet 0
|
||||
OpDecorate %outputValues Binding 1
|
||||
OpMemberDecorate %InputBlock 0 Offset 0
|
||||
OpDecorate %InputBlock BufferBlock
|
||||
OpDecorate %inputValues DescriptorSet 0
|
||||
OpDecorate %inputValues Binding 0
|
||||
OpDecorate %spec_constant SpecId 0
|
||||
%void = OpTypeVoid
|
||||
%void_func = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%int_0 = OpConstant %int 0
|
||||
%int_1 = OpConstant %int 1
|
||||
%int_200 = OpConstant %int 200
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_10 = OpConstant %uint 10
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%Pair = OpTypeStruct %int %int
|
||||
%_ptr_Function_Pair = OpTypePointer Function %Pair
|
||||
%add_pair_members_func_type = OpTypeFunction %int %_ptr_Function_int %_ptr_Function_Pair %_ptr_Function_Pair
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%_arr_int_uint_10 = OpTypeArray %int %uint_10
|
||||
%OutputBlock = OpTypeStruct %_arr_int_uint_10
|
||||
%_ptr_Uniform_OutputBlock = OpTypePointer Uniform %OutputBlock
|
||||
%outputValues = OpVariable %_ptr_Uniform_OutputBlock Uniform
|
||||
%InputBlock = OpTypeStruct %_arr_int_uint_10
|
||||
%_ptr_Uniform_InputBlock = OpTypePointer Uniform %InputBlock
|
||||
%inputValues = OpVariable %_ptr_Uniform_InputBlock Uniform
|
||||
; Replaced %int_100 with an undefined int.
|
||||
%undef_int = OpUndef %int
|
||||
; Composed a spec constant Pair with an undefined int in the second member.
|
||||
%spec_constant = OpSpecConstant %int 0
|
||||
%spec_const_Pair = OpSpecConstantComposite %Pair %spec_constant %undef_int
|
||||
; Composed a constant Pair with the undefined int in the first member.
|
||||
%const_Pair = OpConstantComposite %Pair %undef_int %int_200
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||
%main = OpFunction %void None %void_func
|
||||
%main_label = OpLabel
|
||||
%param_1 = OpVariable %_ptr_Function_int Function
|
||||
%param_2 = OpVariable %_ptr_Function_Pair Function
|
||||
%param_3 = OpVariable %_ptr_Function_Pair Function
|
||||
%gidx_ptr = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||
%gidx = OpLoad %uint %gidx_ptr
|
||||
%input_value_ptr = OpAccessChain %_ptr_Uniform_int %inputValues %int_0 %gidx
|
||||
%input_value = OpLoad %int %input_value_ptr
|
||||
OpStore %param_1 %input_value
|
||||
OpStore %param_2 %spec_const_Pair
|
||||
OpStore %param_3 %const_Pair
|
||||
; Pass the input value as the first argument.
|
||||
; Pass the specialization constant Pair as the second argument.
|
||||
; Pass the constant Pair as the third argument.
|
||||
%retval = OpFunctionCall %int %add_pair_members %param_1 %param_2 %param_3
|
||||
%output_value_ptr = OpAccessChain %_ptr_Uniform_int %outputValues %int_0 %gidx
|
||||
OpStore %output_value_ptr %retval
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%add_pair_members = OpFunction %int None %add_pair_members_func_type
|
||||
%value_ptr = OpFunctionParameter %_ptr_Function_int
|
||||
%pair_1 = OpFunctionParameter %_ptr_Function_Pair
|
||||
%pair_2 = OpFunctionParameter %_ptr_Function_Pair
|
||||
%add_pair_members_label = OpLabel
|
||||
%value = OpLoad %int %value_ptr
|
||||
; Access the first struct member from the first pair.
|
||||
; Access the second struct member from the second pair.
|
||||
; Both should be defined according to the function call above.
|
||||
%pair_1_first_ptr = OpAccessChain %_ptr_Function_int %pair_1 %int_0
|
||||
%pair_2_second_ptr = OpAccessChain %_ptr_Function_int %pair_2 %int_1
|
||||
%pair_1_first = OpLoad %int %pair_1_first_ptr
|
||||
%pair_2_second = OpLoad %int %pair_2_second_ptr
|
||||
%partial_result = OpIAdd %int %value %pair_1_first
|
||||
%final_result = OpIAdd %int %partial_result %pair_2_second
|
||||
OpReturnValue %final_result
|
||||
OpFunctionEnd
|
||||
|
||||
102
shaders/asm/comp/undefined-constant-composite.asm.comp
Normal file
102
shaders/asm/comp/undefined-constant-composite.asm.comp
Normal file
@ -0,0 +1,102 @@
|
||||
;
|
||||
; The shader below is based on the following GLSL shader:
|
||||
;
|
||||
; #version 450
|
||||
;
|
||||
; struct Pair {
|
||||
; int first;
|
||||
; int second;
|
||||
; };
|
||||
;
|
||||
; const Pair constant_pair = { 100, 200 };
|
||||
;
|
||||
; layout(set=0, binding=0, std430) buffer InputBlock {
|
||||
; int array[10];
|
||||
; } inputValues;
|
||||
;
|
||||
; layout(set=0, binding=1, std430) buffer OutputBlock {
|
||||
; int array[10];
|
||||
; } outputValues;
|
||||
;
|
||||
; int add_second (int value, Pair pair) {
|
||||
; return value + pair.second;
|
||||
; }
|
||||
;
|
||||
; void main() {
|
||||
; uint idx = gl_GlobalInvocationID.x;
|
||||
; outputValues.array[idx] = add_second(inputValues.array[idx], constant_pair);
|
||||
; }
|
||||
;
|
||||
; However, the first element of constant_pair has been modified to be undefined.
|
||||
;
|
||||
OpCapability Shader
|
||||
%std450 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID
|
||||
OpExecutionMode %main LocalSize 1 1 1
|
||||
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
|
||||
OpDecorate %_arr_int_uint_10 ArrayStride 4
|
||||
OpMemberDecorate %OutputBlock 0 Offset 0
|
||||
OpDecorate %OutputBlock BufferBlock
|
||||
OpDecorate %outputValues DescriptorSet 0
|
||||
OpDecorate %outputValues Binding 1
|
||||
OpMemberDecorate %InputBlock 0 Offset 0
|
||||
OpDecorate %InputBlock BufferBlock
|
||||
OpDecorate %inputValues DescriptorSet 0
|
||||
OpDecorate %inputValues Binding 0
|
||||
%void = OpTypeVoid
|
||||
%void_func = OpTypeFunction %void
|
||||
%int = OpTypeInt 32 1
|
||||
%uint = OpTypeInt 32 0
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%int_0 = OpConstant %int 0
|
||||
%int_1 = OpConstant %int 1
|
||||
%int_200 = OpConstant %int 200
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_10 = OpConstant %uint 10
|
||||
%_ptr_Function_int = OpTypePointer Function %int
|
||||
%Pair = OpTypeStruct %int %int
|
||||
%_ptr_Function_Pair = OpTypePointer Function %Pair
|
||||
%add_second_func_type = OpTypeFunction %int %_ptr_Function_int %_ptr_Function_Pair
|
||||
%_ptr_Function_uint = OpTypePointer Function %uint
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%_arr_int_uint_10 = OpTypeArray %int %uint_10
|
||||
%OutputBlock = OpTypeStruct %_arr_int_uint_10
|
||||
%_ptr_Uniform_OutputBlock = OpTypePointer Uniform %OutputBlock
|
||||
%outputValues = OpVariable %_ptr_Uniform_OutputBlock Uniform
|
||||
%InputBlock = OpTypeStruct %_arr_int_uint_10
|
||||
%_ptr_Uniform_InputBlock = OpTypePointer Uniform %InputBlock
|
||||
%inputValues = OpVariable %_ptr_Uniform_InputBlock Uniform
|
||||
; Replaced %int_100 with an undefined int.
|
||||
%undef_int = OpUndef %int
|
||||
; Composed a constant Pair with the undefined int in the first member.
|
||||
%const_Pair = OpConstantComposite %Pair %undef_int %int_200
|
||||
%_ptr_Uniform_int = OpTypePointer Uniform %int
|
||||
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||
%main = OpFunction %void None %void_func
|
||||
%main_label = OpLabel
|
||||
%param_1 = OpVariable %_ptr_Function_int Function
|
||||
%param_2 = OpVariable %_ptr_Function_Pair Function
|
||||
%gidx_ptr = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||
%gidx = OpLoad %uint %gidx_ptr
|
||||
%input_value_ptr = OpAccessChain %_ptr_Uniform_int %inputValues %int_0 %gidx
|
||||
%input_value = OpLoad %int %input_value_ptr
|
||||
OpStore %param_1 %input_value
|
||||
OpStore %param_2 %const_Pair
|
||||
%retval = OpFunctionCall %int %add_second %param_1 %param_2
|
||||
%output_value_ptr = OpAccessChain %_ptr_Uniform_int %outputValues %int_0 %gidx
|
||||
OpStore %output_value_ptr %retval
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%add_second = OpFunction %int None %add_second_func_type
|
||||
%value_ptr = OpFunctionParameter %_ptr_Function_int
|
||||
%pair = OpFunctionParameter %_ptr_Function_Pair
|
||||
%add_second_label = OpLabel
|
||||
%value = OpLoad %int %value_ptr
|
||||
; Access the second struct member, which is defined.
|
||||
%pair_second_ptr = OpAccessChain %_ptr_Function_int %pair %int_1
|
||||
%pair_second = OpLoad %int %pair_second_ptr
|
||||
%add_result = OpIAdd %int %value %pair_second
|
||||
OpReturnValue %add_result
|
||||
OpFunctionEnd
|
||||
@ -274,8 +274,6 @@ void CompilerCPP::emit_resources()
|
||||
if (emitted)
|
||||
statement("");
|
||||
|
||||
declare_undefined_values();
|
||||
|
||||
statement("inline void init(spirv_cross_shader& s)");
|
||||
begin_scope();
|
||||
statement(resource_type, "::init(s);");
|
||||
|
||||
@ -66,7 +66,7 @@ ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
|
||||
meta = std::move(other.meta);
|
||||
for (int i = 0; i < TypeCount; i++)
|
||||
ids_for_type[i] = std::move(other.ids_for_type[i]);
|
||||
ids_for_constant_or_type = std::move(other.ids_for_constant_or_type);
|
||||
ids_for_constant_undef_or_type = std::move(other.ids_for_constant_undef_or_type);
|
||||
ids_for_constant_or_variable = std::move(other.ids_for_constant_or_variable);
|
||||
declared_capabilities = std::move(other.declared_capabilities);
|
||||
declared_extensions = std::move(other.declared_extensions);
|
||||
@ -102,7 +102,7 @@ ParsedIR &ParsedIR::operator=(const ParsedIR &other)
|
||||
meta = other.meta;
|
||||
for (int i = 0; i < TypeCount; i++)
|
||||
ids_for_type[i] = other.ids_for_type[i];
|
||||
ids_for_constant_or_type = other.ids_for_constant_or_type;
|
||||
ids_for_constant_undef_or_type = other.ids_for_constant_undef_or_type;
|
||||
ids_for_constant_or_variable = other.ids_for_constant_or_variable;
|
||||
declared_capabilities = other.declared_capabilities;
|
||||
declared_extensions = other.declared_extensions;
|
||||
@ -934,7 +934,7 @@ void ParsedIR::add_typed_id(Types type, ID id)
|
||||
{
|
||||
case TypeConstant:
|
||||
ids_for_constant_or_variable.push_back(id);
|
||||
ids_for_constant_or_type.push_back(id);
|
||||
ids_for_constant_undef_or_type.push_back(id);
|
||||
break;
|
||||
|
||||
case TypeVariable:
|
||||
@ -943,7 +943,8 @@ void ParsedIR::add_typed_id(Types type, ID id)
|
||||
|
||||
case TypeType:
|
||||
case TypeConstantOp:
|
||||
ids_for_constant_or_type.push_back(id);
|
||||
case TypeUndef:
|
||||
ids_for_constant_undef_or_type.push_back(id);
|
||||
break;
|
||||
|
||||
default:
|
||||
|
||||
@ -74,8 +74,8 @@ public:
|
||||
// Special purpose lists which contain a union of types.
|
||||
// This is needed so we can declare specialization constants and structs in an interleaved fashion,
|
||||
// among other things.
|
||||
// Constants can be of struct type, and struct array sizes can use specialization constants.
|
||||
SmallVector<ID> ids_for_constant_or_type;
|
||||
// Constants can be undef or of struct type, and struct array sizes can use specialization constants.
|
||||
SmallVector<ID> ids_for_constant_undef_or_type;
|
||||
SmallVector<ID> ids_for_constant_or_variable;
|
||||
|
||||
// We need to keep track of the width the Ops that contains a type for the
|
||||
|
||||
@ -3416,27 +3416,6 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo
|
||||
statement("");
|
||||
}
|
||||
|
||||
void CompilerGLSL::declare_undefined_values()
|
||||
{
|
||||
bool emitted = false;
|
||||
ir.for_each_typed_id<SPIRUndef>([&](uint32_t, const SPIRUndef &undef) {
|
||||
auto &type = this->get<SPIRType>(undef.basetype);
|
||||
// OpUndef can be void for some reason ...
|
||||
if (type.basetype == SPIRType::Void)
|
||||
return;
|
||||
|
||||
string initializer;
|
||||
if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
|
||||
initializer = join(" = ", to_zero_initialized_expression(undef.basetype));
|
||||
|
||||
statement(variable_decl(type, to_name(undef.self), undef.self), initializer, ";");
|
||||
emitted = true;
|
||||
});
|
||||
|
||||
if (emitted)
|
||||
statement("");
|
||||
}
|
||||
|
||||
bool CompilerGLSL::variable_is_lut(const SPIRVariable &var) const
|
||||
{
|
||||
bool statically_assigned = var.statically_assigned && var.static_expression != ID(0) && var.remapped_variable;
|
||||
@ -3537,7 +3516,7 @@ void CompilerGLSL::emit_resources()
|
||||
//
|
||||
{
|
||||
auto loop_lock = ir.create_loop_hard_lock();
|
||||
for (auto &id_ : ir.ids_for_constant_or_type)
|
||||
for (auto &id_ : ir.ids_for_constant_undef_or_type)
|
||||
{
|
||||
auto &id = ir.ids[id_];
|
||||
|
||||
@ -3590,6 +3569,22 @@ void CompilerGLSL::emit_resources()
|
||||
emit_struct(*type);
|
||||
}
|
||||
}
|
||||
else if (id.get_type() == TypeUndef)
|
||||
{
|
||||
auto &undef = id.get<SPIRUndef>();
|
||||
auto &type = this->get<SPIRType>(undef.basetype);
|
||||
// OpUndef can be void for some reason ...
|
||||
if (type.basetype == SPIRType::Void)
|
||||
return;
|
||||
|
||||
string initializer;
|
||||
if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
|
||||
initializer = join(" = ", to_zero_initialized_expression(undef.basetype));
|
||||
|
||||
// FIXME: If used in a constant, we must declare it as one.
|
||||
statement(variable_decl(type, to_name(undef.self), undef.self), initializer, ";");
|
||||
emitted = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@ -3806,8 +3801,6 @@ void CompilerGLSL::emit_resources()
|
||||
|
||||
if (emitted)
|
||||
statement("");
|
||||
|
||||
declare_undefined_values();
|
||||
}
|
||||
|
||||
void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var)
|
||||
@ -5342,6 +5335,10 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c, bool inside_bloc
|
||||
{
|
||||
res += constant_op_expression(*op);
|
||||
}
|
||||
else if (maybe_get<SPIRUndef>(elem) != nullptr)
|
||||
{
|
||||
res += to_name(elem);
|
||||
}
|
||||
else
|
||||
{
|
||||
auto &subc = get<SPIRConstant>(elem);
|
||||
@ -10064,9 +10061,32 @@ bool CompilerGLSL::should_dereference(uint32_t id)
|
||||
if (auto *var = maybe_get<SPIRVariable>(id))
|
||||
return var->phi_variable;
|
||||
|
||||
// If id is an access chain, we should not dereference it.
|
||||
if (auto *expr = maybe_get<SPIRExpression>(id))
|
||||
return !expr->access_chain;
|
||||
{
|
||||
// If id is an access chain, we should not dereference it.
|
||||
if (expr->access_chain)
|
||||
return false;
|
||||
|
||||
// If id is a forwarded copy of a variable pointer, we should not dereference it.
|
||||
SPIRVariable *var = nullptr;
|
||||
while (expr->loaded_from && expression_is_forwarded(expr->self))
|
||||
{
|
||||
auto &src_type = expression_type(expr->loaded_from);
|
||||
// To be a copy, the pointer and its source expression must be the
|
||||
// same type. Can't check type.self, because for some reason that's
|
||||
// usually the base type with pointers stripped off. This check is
|
||||
// complex enough that I've hoisted it out of the while condition.
|
||||
if (src_type.pointer != type.pointer || src_type.pointer_depth != type.pointer ||
|
||||
src_type.parent_type != type.parent_type)
|
||||
break;
|
||||
if ((var = maybe_get<SPIRVariable>(expr->loaded_from)))
|
||||
break;
|
||||
if (!(expr = maybe_get<SPIRExpression>(expr->loaded_from)))
|
||||
break;
|
||||
}
|
||||
|
||||
return !var || var->phi_variable;
|
||||
}
|
||||
|
||||
// Otherwise, we should dereference this pointer expression.
|
||||
return true;
|
||||
@ -11663,7 +11683,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
|
||||
// RHS expression is immutable, so just forward it.
|
||||
// Copying these things really make no sense, but
|
||||
// seems to be allowed anyways.
|
||||
auto &e = set<SPIRExpression>(id, to_expression(rhs), result_type, true);
|
||||
auto &e = emit_op(result_type, id, to_expression(rhs), true, true);
|
||||
if (pointer)
|
||||
{
|
||||
auto *var = maybe_get_backing_variable(rhs);
|
||||
@ -16900,7 +16920,7 @@ void CompilerGLSL::reorder_type_alias()
|
||||
if (alias_itr < master_itr)
|
||||
{
|
||||
// Must also swap the type order for the constant-type joined array.
|
||||
auto &joined_types = ir.ids_for_constant_or_type;
|
||||
auto &joined_types = ir.ids_for_constant_undef_or_type;
|
||||
auto alt_alias_itr = find(begin(joined_types), end(joined_types), *alias_itr);
|
||||
auto alt_master_itr = find(begin(joined_types), end(joined_types), *master_itr);
|
||||
assert(alt_alias_itr != end(joined_types));
|
||||
|
||||
@ -937,8 +937,6 @@ protected:
|
||||
|
||||
bool type_is_empty(const SPIRType &type);
|
||||
|
||||
virtual void declare_undefined_values();
|
||||
|
||||
bool can_use_io_location(spv::StorageClass storage, bool block);
|
||||
const Instruction *get_next_instruction_in_block(const Instruction &instr);
|
||||
static uint32_t mask_relevant_memory_semantics(uint32_t semantics);
|
||||
|
||||
@ -1409,7 +1409,7 @@ void CompilerHLSL::emit_specialization_constants_and_structs()
|
||||
});
|
||||
|
||||
auto loop_lock = ir.create_loop_hard_lock();
|
||||
for (auto &id_ : ir.ids_for_constant_or_type)
|
||||
for (auto &id_ : ir.ids_for_constant_undef_or_type)
|
||||
{
|
||||
auto &id = ir.ids[id_];
|
||||
|
||||
@ -1471,6 +1471,21 @@ void CompilerHLSL::emit_specialization_constants_and_structs()
|
||||
emit_struct(type);
|
||||
}
|
||||
}
|
||||
else if (id.get_type() == TypeUndef)
|
||||
{
|
||||
auto &undef = id.get<SPIRUndef>();
|
||||
auto &type = this->get<SPIRType>(undef.basetype);
|
||||
// OpUndef can be void for some reason ...
|
||||
if (type.basetype == SPIRType::Void)
|
||||
return;
|
||||
|
||||
string initializer;
|
||||
if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
|
||||
initializer = join(" = ", to_zero_initialized_expression(undef.basetype));
|
||||
|
||||
statement("static ", variable_decl(type, to_name(undef.self), undef.self), initializer, ";");
|
||||
emitted = true;
|
||||
}
|
||||
}
|
||||
|
||||
if (emitted)
|
||||
@ -1514,27 +1529,6 @@ void CompilerHLSL::replace_illegal_names()
|
||||
CompilerGLSL::replace_illegal_names();
|
||||
}
|
||||
|
||||
void CompilerHLSL::declare_undefined_values()
|
||||
{
|
||||
bool emitted = false;
|
||||
ir.for_each_typed_id<SPIRUndef>([&](uint32_t, const SPIRUndef &undef) {
|
||||
auto &type = this->get<SPIRType>(undef.basetype);
|
||||
// OpUndef can be void for some reason ...
|
||||
if (type.basetype == SPIRType::Void)
|
||||
return;
|
||||
|
||||
string initializer;
|
||||
if (options.force_zero_initialized_variables && type_can_zero_initialize(type))
|
||||
initializer = join(" = ", to_zero_initialized_expression(undef.basetype));
|
||||
|
||||
statement("static ", variable_decl(type, to_name(undef.self), undef.self), initializer, ";");
|
||||
emitted = true;
|
||||
});
|
||||
|
||||
if (emitted)
|
||||
statement("");
|
||||
}
|
||||
|
||||
void CompilerHLSL::emit_resources()
|
||||
{
|
||||
auto &execution = get_entry_point();
|
||||
@ -1840,8 +1834,6 @@ void CompilerHLSL::emit_resources()
|
||||
if (emitted)
|
||||
statement("");
|
||||
|
||||
declare_undefined_values();
|
||||
|
||||
if (requires_op_fmod)
|
||||
{
|
||||
static const char *types[] = {
|
||||
|
||||
@ -230,7 +230,6 @@ private:
|
||||
void emit_hlsl_entry_point();
|
||||
void emit_header() override;
|
||||
void emit_resources();
|
||||
void declare_undefined_values() override;
|
||||
void emit_interface_block_globally(const SPIRVariable &type);
|
||||
void emit_interface_block_in_struct(const SPIRVariable &var, std::unordered_set<uint32_t> &active_locations);
|
||||
void emit_interface_block_member_in_struct(const SPIRVariable &var, uint32_t member_index, uint32_t location,
|
||||
|
||||
@ -7217,28 +7217,6 @@ static string inject_top_level_storage_qualifier(const string &expr, const strin
|
||||
}
|
||||
}
|
||||
|
||||
// Undefined global memory is not allowed in MSL.
|
||||
// Declare constant and init to zeros. Use {}, as global constructors can break Metal.
|
||||
void CompilerMSL::declare_undefined_values()
|
||||
{
|
||||
bool emitted = false;
|
||||
ir.for_each_typed_id<SPIRUndef>([&](uint32_t, SPIRUndef &undef) {
|
||||
auto &type = this->get<SPIRType>(undef.basetype);
|
||||
// OpUndef can be void for some reason ...
|
||||
if (type.basetype == SPIRType::Void)
|
||||
return;
|
||||
|
||||
statement(inject_top_level_storage_qualifier(
|
||||
variable_decl(type, to_name(undef.self), undef.self),
|
||||
"constant"),
|
||||
" = {};");
|
||||
emitted = true;
|
||||
});
|
||||
|
||||
if (emitted)
|
||||
statement("");
|
||||
}
|
||||
|
||||
void CompilerMSL::declare_constant_arrays()
|
||||
{
|
||||
bool fully_inlined = ir.ids_for_type[TypeFunction].size() == 1;
|
||||
@ -7304,7 +7282,6 @@ void CompilerMSL::declare_complex_constant_arrays()
|
||||
void CompilerMSL::emit_resources()
|
||||
{
|
||||
declare_constant_arrays();
|
||||
declare_undefined_values();
|
||||
|
||||
// Emit the special [[stage_in]] and [[stage_out]] interface blocks which we created.
|
||||
emit_interface_block(stage_out_var_id);
|
||||
@ -7367,7 +7344,7 @@ void CompilerMSL::emit_specialization_constants_and_structs()
|
||||
emitted = false;
|
||||
declared_structs.clear();
|
||||
|
||||
for (auto &id_ : ir.ids_for_constant_or_type)
|
||||
for (auto &id_ : ir.ids_for_constant_undef_or_type)
|
||||
{
|
||||
auto &id = ir.ids[id_];
|
||||
|
||||
@ -7480,6 +7457,21 @@ void CompilerMSL::emit_specialization_constants_and_structs()
|
||||
emit_struct(get<SPIRType>(type_id));
|
||||
}
|
||||
}
|
||||
else if (id.get_type() == TypeUndef)
|
||||
{
|
||||
auto &undef = id.get<SPIRUndef>();
|
||||
auto &type = get<SPIRType>(undef.basetype);
|
||||
// OpUndef can be void for some reason ...
|
||||
if (type.basetype == SPIRType::Void)
|
||||
return;
|
||||
|
||||
// Undefined global memory is not allowed in MSL.
|
||||
// Declare constant and init to zeros. Use {}, as global constructors can break Metal.
|
||||
statement(
|
||||
inject_top_level_storage_qualifier(variable_decl(type, to_name(undef.self), undef.self), "constant"),
|
||||
" = {};");
|
||||
emitted = true;
|
||||
}
|
||||
}
|
||||
|
||||
if (emitted)
|
||||
@ -9088,12 +9080,33 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
|
||||
uint32_t op0 = ops[2];
|
||||
uint32_t op1 = ops[3];
|
||||
auto &type = get<SPIRType>(result_type);
|
||||
auto input_type = opcode == OpSMulExtended ? int_type : uint_type;
|
||||
auto &output_type = get_type(result_type);
|
||||
string cast_op0, cast_op1;
|
||||
|
||||
auto expected_type = binary_op_bitcast_helper(cast_op0, cast_op1, input_type, op0, op1, false);
|
||||
|
||||
emit_uninitialized_temporary_expression(result_type, result_id);
|
||||
|
||||
statement(to_expression(result_id), ".", to_member_name(type, 0), " = ",
|
||||
to_enclosed_unpacked_expression(op0), " * ", to_enclosed_unpacked_expression(op1), ";");
|
||||
statement(to_expression(result_id), ".", to_member_name(type, 1), " = mulhi(",
|
||||
to_unpacked_expression(op0), ", ", to_unpacked_expression(op1), ");");
|
||||
string mullo_expr, mulhi_expr;
|
||||
mullo_expr = join(cast_op0, " * ", cast_op1);
|
||||
mulhi_expr = join("mulhi(", cast_op0, ", ", cast_op1, ")");
|
||||
|
||||
auto &low_type = get_type(output_type.member_types[0]);
|
||||
auto &high_type = get_type(output_type.member_types[1]);
|
||||
if (low_type.basetype != input_type)
|
||||
{
|
||||
expected_type.basetype = input_type;
|
||||
mullo_expr = join(bitcast_glsl_op(low_type, expected_type), "(", mullo_expr, ")");
|
||||
}
|
||||
if (high_type.basetype != input_type)
|
||||
{
|
||||
expected_type.basetype = input_type;
|
||||
mulhi_expr = join(bitcast_glsl_op(high_type, expected_type), "(", mulhi_expr, ")");
|
||||
}
|
||||
|
||||
statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", mullo_expr, ";");
|
||||
statement(to_expression(result_id), ".", to_member_name(type, 1), " = ", mulhi_expr, ";");
|
||||
break;
|
||||
}
|
||||
|
||||
|
||||
@ -834,7 +834,6 @@ protected:
|
||||
std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override;
|
||||
std::string to_qualifiers_glsl(uint32_t id) override;
|
||||
void replace_illegal_names() override;
|
||||
void declare_undefined_values() override;
|
||||
void declare_constant_arrays();
|
||||
|
||||
void replace_illegal_entry_point_names();
|
||||
|
||||
Loading…
Reference in New Issue
Block a user