Merge pull request #2050 from cdavis5e/op-spec-constant-op-composite-insert
MSL: Implement `CompositeInsert` `OpSpecConstantOp`.
This commit is contained in:
Hans-Kristian Arntzen
GitHub
commit
744279ec78
@ -0,0 +1,111 @@
|
|||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-braces"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
template<typename T, size_t Num>
|
||||||
|
struct spvUnsafeArray
|
||||||
|
{
|
||||||
|
T elements[Num ? Num : 1];
|
||||||
|
|
||||||
|
thread T& operator [] (size_t pos) thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const thread T& operator [] (size_t pos) const thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
device T& operator [] (size_t pos) device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const device T& operator [] (size_t pos) const device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
constexpr const constant T& operator [] (size_t pos) const constant
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
threadgroup T& operator [] (size_t pos) threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
struct _29
|
||||||
|
{
|
||||||
|
spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _m0;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct _7
|
||||||
|
{
|
||||||
|
int _m0[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
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)]];
|
||||||
|
constant int _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 0;
|
||||||
|
constant int _5_tmp [[function_constant(2)]];
|
||||||
|
constant int _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 0;
|
||||||
|
constant spvUnsafeArray<int, 3> _36 = spvUnsafeArray<int, 3>({ _3, 0, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _37 = spvUnsafeArray<int, 3>({ _3, _4, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _38 = spvUnsafeArray<int, 3>({ _3, _4, _5 });
|
||||||
|
constant spvUnsafeArray<int, 3> _39 = spvUnsafeArray<int, 3>({ _4, 0, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _40 = spvUnsafeArray<int, 3>({ _4, _5, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _41 = spvUnsafeArray<int, 3>({ _4, _5, _3 });
|
||||||
|
constant spvUnsafeArray<int, 3> _42 = spvUnsafeArray<int, 3>({ _5, 0, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _43 = spvUnsafeArray<int, 3>({ _5, _3, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _44 = spvUnsafeArray<int, 3>({ _5, _3, _4 });
|
||||||
|
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _45 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
|
||||||
|
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _46 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
|
||||||
|
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _47 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) });
|
||||||
|
constant _29 _48 = _29{ spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) }) };
|
||||||
|
constant int _50 = _48._m0[0][0];
|
||||||
|
constant int _51 = _48._m0[1][0];
|
||||||
|
constant int _52 = _48._m0[0][1];
|
||||||
|
constant int _53 = _48._m0[2][2];
|
||||||
|
constant int _54 = _48._m0[2][0];
|
||||||
|
constant int _55 = _48._m0[1][1];
|
||||||
|
constant bool _56 = (_50 == _51);
|
||||||
|
constant bool _57 = (_52 == _53);
|
||||||
|
constant bool _58 = (_54 == _55);
|
||||||
|
constant int _59 = int(_56);
|
||||||
|
constant int _60 = int(_57);
|
||||||
|
constant int _61 = _58 ? 2 : 1;
|
||||||
|
constant int3 _62 = int3(_3, 0, 0);
|
||||||
|
constant int3 _63 = int3(0, _4, 0);
|
||||||
|
constant int3 _64 = int3(0, 0, _5);
|
||||||
|
constant int3 _65 = int3(_62.x, 0, _62.z);
|
||||||
|
constant int3 _66 = int3(0, _63.y, _63.x);
|
||||||
|
constant int3 _67 = int3(_64.z, 0, _64.z);
|
||||||
|
constant int3 _68 = int3(_65.y, _65.x, _66.y);
|
||||||
|
constant int3 _69 = int3(_67.z, _68.y, _68.z);
|
||||||
|
constant int _70 = _69.x;
|
||||||
|
constant int _71 = _69.y;
|
||||||
|
constant int _72 = _69.z;
|
||||||
|
constant int _73 = (_70 - _71);
|
||||||
|
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,112 @@
|
|||||||
|
#pragma clang diagnostic ignored "-Wmissing-prototypes"
|
||||||
|
#pragma clang diagnostic ignored "-Wmissing-braces"
|
||||||
|
|
||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
using namespace metal;
|
||||||
|
|
||||||
|
template<typename T, size_t Num>
|
||||||
|
struct spvUnsafeArray
|
||||||
|
{
|
||||||
|
T elements[Num ? Num : 1];
|
||||||
|
|
||||||
|
thread T& operator [] (size_t pos) thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const thread T& operator [] (size_t pos) const thread
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
device T& operator [] (size_t pos) device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const device T& operator [] (size_t pos) const device
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
constexpr const constant T& operator [] (size_t pos) const constant
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
|
||||||
|
threadgroup T& operator [] (size_t pos) threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
constexpr const threadgroup T& operator [] (size_t pos) const threadgroup
|
||||||
|
{
|
||||||
|
return elements[pos];
|
||||||
|
}
|
||||||
|
};
|
||||||
|
|
||||||
|
struct _29
|
||||||
|
{
|
||||||
|
spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _m0;
|
||||||
|
};
|
||||||
|
|
||||||
|
struct _7
|
||||||
|
{
|
||||||
|
int _m0[1];
|
||||||
|
};
|
||||||
|
|
||||||
|
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)]];
|
||||||
|
constant int _4 = is_function_constant_defined(_4_tmp) ? _4_tmp : 0;
|
||||||
|
constant int _5_tmp [[function_constant(2)]];
|
||||||
|
constant int _5 = is_function_constant_defined(_5_tmp) ? _5_tmp : 0;
|
||||||
|
constant spvUnsafeArray<int, 3> _36 = spvUnsafeArray<int, 3>({ _3, 0, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _37 = spvUnsafeArray<int, 3>({ _3, _4, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _38 = spvUnsafeArray<int, 3>({ _3, _4, _5 });
|
||||||
|
constant spvUnsafeArray<int, 3> _39 = spvUnsafeArray<int, 3>({ _4, 0, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _40 = spvUnsafeArray<int, 3>({ _4, _5, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _41 = spvUnsafeArray<int, 3>({ _4, _5, _3 });
|
||||||
|
constant spvUnsafeArray<int, 3> _42 = spvUnsafeArray<int, 3>({ _5, 0, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _43 = spvUnsafeArray<int, 3>({ _5, _3, 0 });
|
||||||
|
constant spvUnsafeArray<int, 3> _44 = spvUnsafeArray<int, 3>({ _5, _3, _4 });
|
||||||
|
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _45 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
|
||||||
|
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _46 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ 0, 0, 0 }) });
|
||||||
|
constant spvUnsafeArray<spvUnsafeArray<int, 3>, 3> _47 = spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) });
|
||||||
|
constant _29 _48 = _29{ spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _3 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) }) };
|
||||||
|
constant _29 _49 = _29{ spvUnsafeArray<spvUnsafeArray<int, 3>, 3>({ spvUnsafeArray<int, 3>({ _3, _4, _5 }), spvUnsafeArray<int, 3>({ _4, _5, _5 }), spvUnsafeArray<int, 3>({ _5, _3, _4 }) }) };
|
||||||
|
constant int _50 = _48._m0[0][0];
|
||||||
|
constant int _51 = _48._m0[1][0];
|
||||||
|
constant int _52 = _48._m0[0][1];
|
||||||
|
constant int _53 = _48._m0[2][2];
|
||||||
|
constant int _54 = _48._m0[2][0];
|
||||||
|
constant int _55 = _48._m0[1][1];
|
||||||
|
constant bool _56 = (_50 == _51);
|
||||||
|
constant bool _57 = (_52 == _53);
|
||||||
|
constant bool _58 = (_54 == _55);
|
||||||
|
constant int _59 = int(_56);
|
||||||
|
constant int _60 = int(_57);
|
||||||
|
constant int _61 = _58 ? 2 : 1;
|
||||||
|
constant int3 _62 = int3(_3, 0, 0);
|
||||||
|
constant int3 _63 = int3(0, _4, 0);
|
||||||
|
constant int3 _64 = int3(0, 0, _5);
|
||||||
|
constant int3 _65 = int3(_62.x, 0, _62.z);
|
||||||
|
constant int3 _66 = int3(0, _63.y, _63.x);
|
||||||
|
constant int3 _67 = int3(_64.z, 0, _64.z);
|
||||||
|
constant int3 _68 = int3(_65.y, _65.x, _66.y);
|
||||||
|
constant int3 _69 = int3(_67.z, _68.y, _68.z);
|
||||||
|
constant int _70 = _69.x;
|
||||||
|
constant int _71 = _69.y;
|
||||||
|
constant int _72 = _69.z;
|
||||||
|
constant int _73 = (_70 - _71);
|
||||||
|
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);
|
||||||
|
}
|
||||||
|
|
||||||
107
shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp
Normal file
107
shaders-msl/asm/comp/op-spec-constant-op-vector-related.asm.comp
Normal file
@ -0,0 +1,107 @@
|
|||||||
|
OpCapability Shader
|
||||||
|
OpMemoryModel Logical GLSL450
|
||||||
|
OpEntryPoint GLCompute %main "main" %id
|
||||||
|
OpExecutionMode %main LocalSize 1 1 1
|
||||||
|
OpName %main "main"
|
||||||
|
OpName %id "gl_GlobalInvocationID"
|
||||||
|
OpDecorate %id BuiltIn GlobalInvocationId
|
||||||
|
OpDecorate %sc_0 SpecId 0
|
||||||
|
OpDecorate %sc_1 SpecId 1
|
||||||
|
OpDecorate %sc_2 SpecId 2
|
||||||
|
OpDecorate %i32arr ArrayStride 4
|
||||||
|
OpDecorate %buf BufferBlock
|
||||||
|
OpDecorate %indata DescriptorSet 0
|
||||||
|
OpDecorate %indata Binding 0
|
||||||
|
OpDecorate %outdata DescriptorSet 0
|
||||||
|
OpDecorate %outdata Binding 1
|
||||||
|
OpDecorate %f32arr ArrayStride 4
|
||||||
|
OpMemberDecorate %buf 0 Offset 0
|
||||||
|
%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
|
||||||
|
%ivec3 = OpTypeVector %i32 3
|
||||||
|
%zero = OpConstant %i32 0
|
||||||
|
%one = OpConstant %i32 1
|
||||||
|
%two = OpConstant %i32 2
|
||||||
|
%three = OpConstant %i32 3
|
||||||
|
%iarr3 = OpTypeArray %i32 %three
|
||||||
|
%imat3 = OpTypeArray %iarr3 %three
|
||||||
|
%struct = OpTypeStruct %imat3
|
||||||
|
%buf = OpTypeStruct %i32arr
|
||||||
|
%bufptr = OpTypePointer Uniform %buf
|
||||||
|
%indata = OpVariable %bufptr Uniform
|
||||||
|
%outdata = OpVariable %bufptr Uniform
|
||||||
|
%id = OpVariable %uvec3ptr Input
|
||||||
|
%ivec3_0 = OpConstantComposite %ivec3 %zero %zero %zero
|
||||||
|
%vec3_undef = OpUndef %ivec3
|
||||||
|
%iarr3_0 = OpConstantComposite %iarr3 %zero %zero %zero
|
||||||
|
%imat3_0 = OpConstantComposite %imat3 %iarr3_0 %iarr3_0 %iarr3_0
|
||||||
|
%struct_0 = OpConstantComposite %struct %imat3_0
|
||||||
|
%sc_0 = OpSpecConstant %i32 0
|
||||||
|
%sc_1 = OpSpecConstant %i32 0
|
||||||
|
%sc_2 = OpSpecConstant %i32 0
|
||||||
|
%iarr3_a = OpSpecConstantOp %iarr3 CompositeInsert %sc_0 %iarr3_0 0
|
||||||
|
%iarr3_b = OpSpecConstantOp %iarr3 CompositeInsert %sc_1 %iarr3_a 1
|
||||||
|
%iarr3_c = OpSpecConstantOp %iarr3 CompositeInsert %sc_2 %iarr3_b 2
|
||||||
|
%iarr3_d = OpSpecConstantOp %iarr3 CompositeInsert %sc_1 %iarr3_0 0
|
||||||
|
%iarr3_e = OpSpecConstantOp %iarr3 CompositeInsert %sc_2 %iarr3_d 1
|
||||||
|
%iarr3_f = OpSpecConstantOp %iarr3 CompositeInsert %sc_0 %iarr3_e 2
|
||||||
|
%iarr3_g = OpSpecConstantOp %iarr3 CompositeInsert %sc_2 %iarr3_0 0
|
||||||
|
%iarr3_h = OpSpecConstantOp %iarr3 CompositeInsert %sc_0 %iarr3_g 1
|
||||||
|
%iarr3_i = OpSpecConstantOp %iarr3 CompositeInsert %sc_1 %iarr3_h 2
|
||||||
|
%imat3_a = OpSpecConstantOp %imat3 CompositeInsert %iarr3_c %imat3_0 0
|
||||||
|
%imat3_b = OpSpecConstantOp %imat3 CompositeInsert %iarr3_f %imat3_a 1
|
||||||
|
%imat3_c = OpSpecConstantOp %imat3 CompositeInsert %iarr3_i %imat3_b 2
|
||||||
|
%struct_a = OpSpecConstantOp %struct CompositeInsert %imat3_c %struct_0 0
|
||||||
|
%struct_b = OpSpecConstantOp %struct CompositeInsert %sc_2 %struct_a 0 1 2
|
||||||
|
%comp_0_0 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 0 0
|
||||||
|
%comp_1_0 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 1 0
|
||||||
|
%comp_0_1 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 0 1
|
||||||
|
%comp_2_2 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 2 2
|
||||||
|
%comp_2_0 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 2 0
|
||||||
|
%comp_1_1 = OpSpecConstantOp %i32 CompositeExtract %struct_a 0 1 1
|
||||||
|
%cmpres_0 = OpSpecConstantOp %bool IEqual %comp_0_0 %comp_1_0
|
||||||
|
%cmpres_1 = OpSpecConstantOp %bool IEqual %comp_0_1 %comp_2_2
|
||||||
|
%cmpres_2 = OpSpecConstantOp %bool IEqual %comp_2_0 %comp_1_1
|
||||||
|
%mustbe_0 = OpSpecConstantOp %i32 Select %cmpres_0 %one %zero
|
||||||
|
%mustbe_1 = OpSpecConstantOp %i32 Select %cmpres_1 %one %zero
|
||||||
|
%mustbe_2 = OpSpecConstantOp %i32 Select %cmpres_2 %two %one
|
||||||
|
%sc_vec3_0 = OpSpecConstantOp %ivec3 CompositeInsert %sc_0 %ivec3_0 0
|
||||||
|
%sc_vec3_1 = OpSpecConstantOp %ivec3 CompositeInsert %sc_1 %ivec3_0 1
|
||||||
|
%sc_vec3_2 = OpSpecConstantOp %ivec3 CompositeInsert %sc_2 %ivec3_0 2
|
||||||
|
%sc_vec3_0_s = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_0 %vec3_undef 0 0xFFFFFFFF 2
|
||||||
|
%sc_vec3_1_s = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_1 %vec3_undef 0xFFFFFFFF 1 0
|
||||||
|
%sc_vec3_2_s = OpSpecConstantOp %ivec3 VectorShuffle %vec3_undef %sc_vec3_2 5 0xFFFFFFFF 5
|
||||||
|
%sc_vec3_01 = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_0_s %sc_vec3_1_s 1 0 4
|
||||||
|
%sc_vec3_012 = OpSpecConstantOp %ivec3 VectorShuffle %sc_vec3_01 %sc_vec3_2_s 5 1 2
|
||||||
|
%sc_ext_0 = OpSpecConstantOp %i32 CompositeExtract %sc_vec3_012 0
|
||||||
|
%sc_ext_1 = OpSpecConstantOp %i32 CompositeExtract %sc_vec3_012 1
|
||||||
|
%sc_ext_2 = OpSpecConstantOp %i32 CompositeExtract %sc_vec3_012 2
|
||||||
|
%sc_sub = OpSpecConstantOp %i32 ISub %sc_ext_0 %sc_ext_1
|
||||||
|
%sc_factor = OpSpecConstantOp %i32 IMul %sc_sub %sc_ext_2
|
||||||
|
%main = OpFunction %void None %voidf
|
||||||
|
%label = OpLabel
|
||||||
|
%subf_a = OpISub %i32 %one %mustbe_0
|
||||||
|
%subf_b = OpIMul %i32 %subf_a %mustbe_1
|
||||||
|
%subf_c = OpISub %i32 %mustbe_2 %one
|
||||||
|
%factor = OpIMul %i32 %subf_b %subf_c
|
||||||
|
%sc_final = OpIMul %i32 %factor %sc_factor
|
||||||
|
%idval = OpLoad %uvec3 %id
|
||||||
|
%x = OpCompositeExtract %u32 %idval 0
|
||||||
|
%inloc = OpAccessChain %i32ptr %indata %zero %x
|
||||||
|
%inval = OpLoad %i32 %inloc
|
||||||
|
%final = OpIAdd %i32 %inval %sc_final
|
||||||
|
%outloc = OpAccessChain %i32ptr %outdata %zero %x
|
||||||
|
OpStore %outloc %final
|
||||||
|
OpReturn
|
||||||
|
OpFunctionEnd
|
||||||
@ -210,7 +210,8 @@ public:
|
|||||||
buffer_capacity = N;
|
buffer_capacity = N;
|
||||||
}
|
}
|
||||||
|
|
||||||
SmallVector(const T *arg_list_begin, const T *arg_list_end) SPIRV_CROSS_NOEXCEPT : SmallVector()
|
template <typename U>
|
||||||
|
SmallVector(const U *arg_list_begin, const U *arg_list_end) SPIRV_CROSS_NOEXCEPT : SmallVector()
|
||||||
{
|
{
|
||||||
auto count = size_t(arg_list_end - arg_list_begin);
|
auto count = size_t(arg_list_end - arg_list_begin);
|
||||||
reserve(count);
|
reserve(count);
|
||||||
@ -219,7 +220,13 @@ public:
|
|||||||
this->buffer_size = count;
|
this->buffer_size = count;
|
||||||
}
|
}
|
||||||
|
|
||||||
SmallVector(std::initializer_list<T> init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end())
|
template <typename U>
|
||||||
|
SmallVector(std::initializer_list<U> init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end())
|
||||||
|
{
|
||||||
|
}
|
||||||
|
|
||||||
|
template <typename U, size_t M>
|
||||||
|
SmallVector(const U (&init)[M]) SPIRV_CROSS_NOEXCEPT : SmallVector(init, init + M)
|
||||||
{
|
{
|
||||||
}
|
}
|
||||||
|
|
||||||
|
|||||||
142
spirv_glsl.cpp
142
spirv_glsl.cpp
@ -4859,6 +4859,9 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read)
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
if (expression_is_forwarded(id))
|
||||||
|
return constant_expression(c);
|
||||||
|
|
||||||
return to_name(id);
|
return to_name(id);
|
||||||
}
|
}
|
||||||
else if (c.is_used_as_lut)
|
else if (c.is_used_as_lut)
|
||||||
@ -4930,6 +4933,80 @@ string CompilerGLSL::to_expression(uint32_t id, bool register_expression_read)
|
|||||||
}
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
|
SmallVector<ConstantID> CompilerGLSL::get_composite_constant_ids(ConstantID const_id)
|
||||||
|
{
|
||||||
|
if (auto *constant = maybe_get<SPIRConstant>(const_id))
|
||||||
|
{
|
||||||
|
const auto &type = get<SPIRType>(constant->constant_type);
|
||||||
|
if (is_array(type) || type.basetype == SPIRType::Struct)
|
||||||
|
return constant->subconstants;
|
||||||
|
if (is_matrix(type))
|
||||||
|
return constant->m.id;
|
||||||
|
if (is_vector(type))
|
||||||
|
return constant->m.c[0].id;
|
||||||
|
SPIRV_CROSS_THROW("Unexpected scalar constant!");
|
||||||
|
}
|
||||||
|
if (!const_composite_insert_ids.count(const_id))
|
||||||
|
SPIRV_CROSS_THROW("Unimplemented for this OpSpecConstantOp!");
|
||||||
|
return const_composite_insert_ids[const_id];
|
||||||
|
}
|
||||||
|
|
||||||
|
void CompilerGLSL::fill_composite_constant(SPIRConstant &constant, TypeID type_id,
|
||||||
|
const SmallVector<ConstantID> &initializers)
|
||||||
|
{
|
||||||
|
auto &type = get<SPIRType>(type_id);
|
||||||
|
constant.specialization = true;
|
||||||
|
if (is_array(type) || type.basetype == SPIRType::Struct)
|
||||||
|
{
|
||||||
|
constant.subconstants = initializers;
|
||||||
|
}
|
||||||
|
else if (is_matrix(type))
|
||||||
|
{
|
||||||
|
constant.m.columns = type.columns;
|
||||||
|
for (uint32_t i = 0; i < type.columns; ++i)
|
||||||
|
{
|
||||||
|
constant.m.id[i] = initializers[i];
|
||||||
|
constant.m.c[i].vecsize = type.vecsize;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
else if (is_vector(type))
|
||||||
|
{
|
||||||
|
constant.m.c[0].vecsize = type.vecsize;
|
||||||
|
for (uint32_t i = 0; i < type.vecsize; ++i)
|
||||||
|
constant.m.c[0].id[i] = initializers[i];
|
||||||
|
}
|
||||||
|
else
|
||||||
|
SPIRV_CROSS_THROW("Unexpected scalar in SpecConstantOp CompositeInsert!");
|
||||||
|
}
|
||||||
|
|
||||||
|
void CompilerGLSL::set_composite_constant(ConstantID const_id, TypeID type_id,
|
||||||
|
const SmallVector<ConstantID> &initializers)
|
||||||
|
{
|
||||||
|
if (maybe_get<SPIRConstantOp>(const_id))
|
||||||
|
{
|
||||||
|
const_composite_insert_ids[const_id] = initializers;
|
||||||
|
return;
|
||||||
|
}
|
||||||
|
|
||||||
|
auto &constant = set<SPIRConstant>(const_id, type_id);
|
||||||
|
fill_composite_constant(constant, type_id, initializers);
|
||||||
|
forwarded_temporaries.insert(const_id);
|
||||||
|
}
|
||||||
|
|
||||||
|
TypeID CompilerGLSL::get_composite_member_type(TypeID type_id, uint32_t member_idx)
|
||||||
|
{
|
||||||
|
auto &type = get<SPIRType>(type_id);
|
||||||
|
if (is_array(type))
|
||||||
|
return type.parent_type;
|
||||||
|
if (type.basetype == SPIRType::Struct)
|
||||||
|
return type.member_types[member_idx];
|
||||||
|
if (is_matrix(type))
|
||||||
|
return type.parent_type;
|
||||||
|
if (is_vector(type))
|
||||||
|
return type.parent_type;
|
||||||
|
SPIRV_CROSS_THROW("Shouldn't reach lower than vector handling OpSpecConstantOp CompositeInsert!");
|
||||||
|
}
|
||||||
|
|
||||||
string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
||||||
{
|
{
|
||||||
auto &type = get<SPIRType>(cop.basetype);
|
auto &type = get<SPIRType>(cop.basetype);
|
||||||
@ -5034,10 +5111,21 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
|||||||
for (uint32_t i = 2; i < uint32_t(cop.arguments.size()); i++)
|
for (uint32_t i = 2; i < uint32_t(cop.arguments.size()); i++)
|
||||||
{
|
{
|
||||||
uint32_t index = cop.arguments[i];
|
uint32_t index = cop.arguments[i];
|
||||||
if (index >= left_components)
|
if (index == 0xFFFFFFFF)
|
||||||
|
{
|
||||||
|
SPIRConstant c;
|
||||||
|
c.constant_type = type.parent_type;
|
||||||
|
assert(type.parent_type != ID(0));
|
||||||
|
expr += constant_expression(c);
|
||||||
|
}
|
||||||
|
else if (index >= left_components)
|
||||||
|
{
|
||||||
expr += right_arg + "." + "xyzw"[index - left_components];
|
expr += right_arg + "." + "xyzw"[index - left_components];
|
||||||
|
}
|
||||||
else
|
else
|
||||||
|
{
|
||||||
expr += left_arg + "." + "xyzw"[index];
|
expr += left_arg + "." + "xyzw"[index];
|
||||||
|
}
|
||||||
|
|
||||||
if (i + 1 < uint32_t(cop.arguments.size()))
|
if (i + 1 < uint32_t(cop.arguments.size()))
|
||||||
expr += ", ";
|
expr += ", ";
|
||||||
@ -5055,7 +5143,30 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
|
|||||||
}
|
}
|
||||||
|
|
||||||
case OpCompositeInsert:
|
case OpCompositeInsert:
|
||||||
SPIRV_CROSS_THROW("OpCompositeInsert spec constant op is not supported.");
|
{
|
||||||
|
SmallVector<ConstantID> new_init = get_composite_constant_ids(cop.arguments[1]);
|
||||||
|
uint32_t idx;
|
||||||
|
uint32_t target_id = cop.self;
|
||||||
|
uint32_t target_type_id = cop.basetype;
|
||||||
|
// We have to drill down to the part we want to modify, and create new
|
||||||
|
// constants for each containing part.
|
||||||
|
for (idx = 2; idx < cop.arguments.size() - 1; ++idx)
|
||||||
|
{
|
||||||
|
uint32_t new_const = ir.increase_bound_by(1);
|
||||||
|
uint32_t old_const = new_init[cop.arguments[idx]];
|
||||||
|
new_init[cop.arguments[idx]] = new_const;
|
||||||
|
set_composite_constant(target_id, target_type_id, new_init);
|
||||||
|
new_init = get_composite_constant_ids(old_const);
|
||||||
|
target_id = new_const;
|
||||||
|
target_type_id = get_composite_member_type(target_type_id, cop.arguments[idx]);
|
||||||
|
}
|
||||||
|
// Now replace the initializer with the one from this instruction.
|
||||||
|
new_init[cop.arguments[idx]] = cop.arguments[0];
|
||||||
|
set_composite_constant(target_id, target_type_id, new_init);
|
||||||
|
SPIRConstant tmp_const(cop.basetype);
|
||||||
|
fill_composite_constant(tmp_const, cop.basetype, const_composite_insert_ids[cop.self]);
|
||||||
|
return constant_expression(tmp_const);
|
||||||
|
}
|
||||||
|
|
||||||
default:
|
default:
|
||||||
// Some opcodes are unimplemented here, these are currently not possible to test from glslang.
|
// Some opcodes are unimplemented here, these are currently not possible to test from glslang.
|
||||||
@ -5206,20 +5317,27 @@ string CompilerGLSL::constant_expression(const SPIRConstant &c, bool inside_bloc
|
|||||||
uint32_t subconstant_index = 0;
|
uint32_t subconstant_index = 0;
|
||||||
for (auto &elem : c.subconstants)
|
for (auto &elem : c.subconstants)
|
||||||
{
|
{
|
||||||
auto &subc = get<SPIRConstant>(elem);
|
if (auto *op = maybe_get<SPIRConstantOp>(elem))
|
||||||
if (subc.specialization)
|
{
|
||||||
res += to_name(elem);
|
res += constant_op_expression(*op);
|
||||||
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
if (type.array.empty() && type.basetype == SPIRType::Struct)
|
auto &subc = get<SPIRConstant>(elem);
|
||||||
|
if (subc.specialization && !expression_is_forwarded(elem))
|
||||||
|
res += to_name(elem);
|
||||||
|
else
|
||||||
{
|
{
|
||||||
// When we get down to emitting struct members, override the block-like information.
|
if (type.array.empty() && type.basetype == SPIRType::Struct)
|
||||||
// For constants, we can freely mix and match block-like state.
|
{
|
||||||
inside_block_like_struct_scope =
|
// When we get down to emitting struct members, override the block-like information.
|
||||||
has_member_decoration(type.self, subconstant_index, DecorationOffset);
|
// For constants, we can freely mix and match block-like state.
|
||||||
}
|
inside_block_like_struct_scope =
|
||||||
|
has_member_decoration(type.self, subconstant_index, DecorationOffset);
|
||||||
|
}
|
||||||
|
|
||||||
res += constant_expression(subc, inside_block_like_struct_scope);
|
res += constant_expression(subc, inside_block_like_struct_scope);
|
||||||
|
}
|
||||||
}
|
}
|
||||||
|
|
||||||
if (&elem != &c.subconstants.back())
|
if (&elem != &c.subconstants.back())
|
||||||
|
|||||||
@ -989,6 +989,12 @@ protected:
|
|||||||
|
|
||||||
private:
|
private:
|
||||||
void init();
|
void init();
|
||||||
|
|
||||||
|
SmallVector<ConstantID> get_composite_constant_ids(ConstantID const_id);
|
||||||
|
void fill_composite_constant(SPIRConstant &constant, TypeID type_id, const SmallVector<ConstantID> &initializers);
|
||||||
|
void set_composite_constant(ConstantID const_id, TypeID type_id, const SmallVector<ConstantID> &initializers);
|
||||||
|
TypeID get_composite_member_type(TypeID type_id, uint32_t member_idx);
|
||||||
|
std::unordered_map<uint32_t, SmallVector<ConstantID>> const_composite_insert_ids;
|
||||||
};
|
};
|
||||||
} // namespace SPIRV_CROSS_NAMESPACE
|
} // namespace SPIRV_CROSS_NAMESPACE
|
||||||
|
|
||||||
|
|||||||
Loading…
Reference in New Issue
Block a user