Merge pull request #2242 from KhronosGroup/fix-2221
MSL: Improve PtrAccessChain handling.
This commit is contained in:
Hans-Kristian Arntzen
GitHub
commit
9da5f7ce45
@ -0,0 +1,18 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
|
||||
|
||||
kernel void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
threadgroup float2 test[64];
|
||||
float _21 = float(gl_GlobalInvocationID.x);
|
||||
float2 _22 = float2(_21);
|
||||
((&((&test)[0u]))[0u])[1u + 2u] = _22;
|
||||
((&test)[0u])[1u + 2u] = _22;
|
||||
((&test)[0u])[3u] = _22;
|
||||
((threadgroup float*)&((&test)[0u])[2u])[0u + 1u] = _21;
|
||||
}
|
||||
|
||||
@ -0,0 +1,14 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(64u, 1u, 1u);
|
||||
|
||||
kernel void main0(uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
|
||||
{
|
||||
threadgroup float2 test[64];
|
||||
float _21 = float(gl_GlobalInvocationID.x);
|
||||
((threadgroup float*)&(*(true ? &test[1u] : &test[2u])))[1u] = _21;
|
||||
}
|
||||
|
||||
@ -0,0 +1,134 @@
|
||||
#include <metal_stdlib>
|
||||
#include <simd/simd.h>
|
||||
|
||||
using namespace metal;
|
||||
|
||||
constant uint _15_tmp [[function_constant(0)]];
|
||||
constant uint _15 = is_function_constant_defined(_15_tmp) ? _15_tmp : 1u;
|
||||
constant uint _16_tmp [[function_constant(1)]];
|
||||
constant uint _16 = is_function_constant_defined(_16_tmp) ? _16_tmp : 1u;
|
||||
constant uint _17_tmp [[function_constant(2)]];
|
||||
constant uint _17 = is_function_constant_defined(_17_tmp) ? _17_tmp : 1u;
|
||||
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(_15, _16, _17);
|
||||
|
||||
struct _6
|
||||
{
|
||||
uint4 _m0[1];
|
||||
};
|
||||
|
||||
struct _7
|
||||
{
|
||||
uint _m0;
|
||||
};
|
||||
|
||||
struct _8
|
||||
{
|
||||
_7 _m0;
|
||||
};
|
||||
|
||||
constant uchar4 _137 = {};
|
||||
|
||||
kernel void main0(device _6& _25 [[buffer(0)]], constant _8& _29 [[buffer(1)]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]], uint3 gl_WorkGroupID [[threadgroup_position_in_grid]])
|
||||
{
|
||||
threadgroup uint _5[256];
|
||||
threadgroup uchar _10[1024];
|
||||
uint3 _20 = gl_WorkGroupSize;
|
||||
bool _40 = _29._m0._m0 != 0u;
|
||||
if (_40)
|
||||
{
|
||||
uchar _58 = uchar(((gl_LocalInvocationID.y * gl_LocalInvocationID.x) / gl_WorkGroupID.y) % 255u);
|
||||
uint _66;
|
||||
uint _61 = 0u;
|
||||
uint _62;
|
||||
for (;;)
|
||||
{
|
||||
_62 = _61 * _29._m0._m0;
|
||||
_66 = 0u;
|
||||
for (;;)
|
||||
{
|
||||
uint _67 = _66 + _62;
|
||||
uint _68 = _66 * _61;
|
||||
_5[_67] = gl_WorkGroupID.x + _68;
|
||||
uint _74 = _67 << 2u;
|
||||
uint _76 = _74 >> 10u;
|
||||
uint _78 = _74 & 1020u;
|
||||
uchar4 _80 = as_type<uchar4>(gl_WorkGroupID.y + _68);
|
||||
((&_10)[_76])[_78 | 1u] = _80.y;
|
||||
((&_10)[_76])[_78 | 2u] = _80.z;
|
||||
((&_10)[_76])[_78 | 3u] = _80.w;
|
||||
((&_10)[_76])[_78] = _58;
|
||||
uint _93 = _66 + 1u;
|
||||
if (_93 >= _29._m0._m0)
|
||||
{
|
||||
break;
|
||||
}
|
||||
else
|
||||
{
|
||||
_66 = _93;
|
||||
}
|
||||
}
|
||||
uint _100 = _61 + 1u;
|
||||
if (_100 >= _29._m0._m0)
|
||||
{
|
||||
break;
|
||||
}
|
||||
else
|
||||
{
|
||||
_61 = _100;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
uint _112;
|
||||
if (_40)
|
||||
{
|
||||
_112 = 0u;
|
||||
uint _117;
|
||||
uint _113;
|
||||
for (;;)
|
||||
{
|
||||
_113 = _112 * _29._m0._m0;
|
||||
_117 = 0u;
|
||||
for (;;)
|
||||
{
|
||||
uint _118 = _117 + _113;
|
||||
uint _123 = _118 << 2u;
|
||||
uint _124 = _123 >> 10u;
|
||||
uint _125 = _123 & 1020u;
|
||||
uchar4 _138;
|
||||
_138.x = ((&_10)[_124])[_125];
|
||||
_138.y = ((&_10)[_124])[_125 | 1u];
|
||||
_138.z = ((&_10)[_124])[_125 | 2u];
|
||||
_138.w = ((&_10)[_124])[_125 | 3u];
|
||||
uint _143 = _5[_118] + as_type<uint>(_138);
|
||||
uint4 _144 = _25._m0[_118];
|
||||
_144.x = _143;
|
||||
_144.y = _143 >> 2u;
|
||||
_144.w = _143 >> 3u;
|
||||
_25._m0[_118] = _144;
|
||||
uint _150 = _117 + 1u;
|
||||
if (_150 >= _29._m0._m0)
|
||||
{
|
||||
break;
|
||||
}
|
||||
else
|
||||
{
|
||||
_117 = _150;
|
||||
}
|
||||
}
|
||||
uint _157 = _112 + 1u;
|
||||
if (_157 >= _29._m0._m0)
|
||||
{
|
||||
break;
|
||||
}
|
||||
else
|
||||
{
|
||||
_112 = _157;
|
||||
continue;
|
||||
}
|
||||
}
|
||||
}
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
|
||||
71
shaders-msl-no-opt/asm/comp/variable-pointers-2.asm.comp
Normal file
71
shaders-msl-no-opt/asm/comp/variable-pointers-2.asm.comp
Normal file
@ -0,0 +1,71 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 11
|
||||
; Bound: 26
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability VariablePointers
|
||||
OpExtension "SPV_KHR_variable_pointers"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_LocalInvocationIndex %gl_GlobalInvocationID
|
||||
OpExecutionMode %main LocalSize 64 1 1
|
||||
OpSource GLSL 450
|
||||
OpName %main "main"
|
||||
OpName %test "test"
|
||||
OpName %gl_LocalInvocationIndex "gl_LocalInvocationIndex"
|
||||
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
|
||||
OpDecorate %gl_LocalInvocationIndex BuiltIn LocalInvocationIndex
|
||||
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
|
||||
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%v2float = OpTypeVector %float 2
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_64 = OpConstant %uint 64
|
||||
%_arr_v2float_uint_64 = OpTypeArray %v2float %uint_64
|
||||
%_ptr_Workgroup__arr_v2float_uint_64 = OpTypePointer Workgroup %_arr_v2float_uint_64
|
||||
%test = OpVariable %_ptr_Workgroup__arr_v2float_uint_64 Workgroup
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%gl_LocalInvocationIndex = OpVariable %_ptr_Input_uint Input
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_ptr_Workgroup_float = OpTypePointer Workgroup %float
|
||||
%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
|
||||
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_64 %uint_1 %uint_1
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%14 = OpLoad %uint %gl_LocalInvocationIndex
|
||||
%19 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||
%20 = OpLoad %uint %19
|
||||
%21 = OpConvertUToF %float %20
|
||||
%22 = OpCompositeConstruct %v2float %21 %21
|
||||
|
||||
; Dummy expression. *(&test + 0)
|
||||
%ptr0 = OpPtrAccessChain %_ptr_Workgroup__arr_v2float_uint_64 %test %uint_0
|
||||
%ptr1 = OpPtrAccessChain %_ptr_Workgroup_v2float %ptr0 %uint_0 %uint_1
|
||||
%ptr2 = OpPtrAccessChain %_ptr_Workgroup_v2float %ptr1 %uint_2
|
||||
OpStore %ptr2 %22
|
||||
|
||||
; Chain PtrAccessChain while keeping pointer type.
|
||||
%ptr3 = OpPtrAccessChain %_ptr_Workgroup_v2float %test %uint_0 %uint_1
|
||||
%ptr4 = OpPtrAccessChain %_ptr_Workgroup_v2float %ptr3 %uint_2
|
||||
OpStore %ptr4 %22
|
||||
|
||||
; Same semantics.
|
||||
%ptr5 = OpPtrAccessChain %_ptr_Workgroup_v2float %test %uint_0 %uint_3
|
||||
OpStore %ptr5 %22
|
||||
|
||||
; Scalar shenanigans.
|
||||
%ptr6 = OpPtrAccessChain %_ptr_Workgroup_float %test %uint_0 %uint_2 %uint_0
|
||||
%ptr7 = OpPtrAccessChain %_ptr_Workgroup_float %ptr6 %uint_1
|
||||
OpStore %ptr7 %21
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -0,0 +1,60 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Khronos Glslang Reference Front End; 11
|
||||
; Bound: 26
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability VariablePointers
|
||||
OpExtension "SPV_KHR_variable_pointers"
|
||||
%1 = OpExtInstImport "GLSL.std.450"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %main "main" %gl_LocalInvocationIndex %gl_GlobalInvocationID
|
||||
OpExecutionMode %main LocalSize 64 1 1
|
||||
OpSource GLSL 450
|
||||
OpName %main "main"
|
||||
OpName %test "test"
|
||||
OpName %gl_LocalInvocationIndex "gl_LocalInvocationIndex"
|
||||
OpName %gl_GlobalInvocationID "gl_GlobalInvocationID"
|
||||
OpDecorate %gl_LocalInvocationIndex BuiltIn LocalInvocationIndex
|
||||
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
|
||||
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
|
||||
%void = OpTypeVoid
|
||||
%3 = OpTypeFunction %void
|
||||
%float = OpTypeFloat 32
|
||||
%bool = OpTypeBool
|
||||
%true = OpConstantTrue %bool
|
||||
%v2float = OpTypeVector %float 2
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_64 = OpConstant %uint 64
|
||||
%_arr_v2float_uint_64 = OpTypeArray %v2float %uint_64
|
||||
%_ptr_Workgroup__arr_v2float_uint_64 = OpTypePointer Workgroup %_arr_v2float_uint_64
|
||||
%test = OpVariable %_ptr_Workgroup__arr_v2float_uint_64 Workgroup
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%gl_LocalInvocationIndex = OpVariable %_ptr_Input_uint Input
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%_ptr_Workgroup_float = OpTypePointer Workgroup %float
|
||||
%_ptr_Workgroup_v2float = OpTypePointer Workgroup %v2float
|
||||
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_64 %uint_1 %uint_1
|
||||
%main = OpFunction %void None %3
|
||||
%5 = OpLabel
|
||||
%14 = OpLoad %uint %gl_LocalInvocationIndex
|
||||
%19 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0
|
||||
%20 = OpLoad %uint %19
|
||||
%21 = OpConvertUToF %float %20
|
||||
%22 = OpCompositeConstruct %v2float %21 %21
|
||||
|
||||
%a = OpAccessChain %_ptr_Workgroup_v2float %test %uint_1
|
||||
%b = OpAccessChain %_ptr_Workgroup_v2float %test %uint_2
|
||||
%c = OpSelect %_ptr_Workgroup_v2float %true %a %b
|
||||
|
||||
%d = OpAccessChain %_ptr_Workgroup_float %c %uint_1
|
||||
OpStore %d %21
|
||||
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
@ -0,0 +1,214 @@
|
||||
; SPIR-V
|
||||
; Version: 1.0
|
||||
; Generator: Google Clspv; 0
|
||||
; Bound: 175
|
||||
; Schema: 0
|
||||
OpCapability Shader
|
||||
OpCapability Int8
|
||||
OpCapability VariablePointers
|
||||
OpExtension "SPV_KHR_storage_buffer_storage_class"
|
||||
OpExtension "SPV_KHR_variable_pointers"
|
||||
OpExtension "SPV_KHR_non_semantic_info"
|
||||
%163 = OpExtInstImport "NonSemantic.ClspvReflection.5"
|
||||
OpMemoryModel Logical GLSL450
|
||||
OpEntryPoint GLCompute %32 "main" %gl_LocalInvocationID %gl_WorkGroupID
|
||||
OpSource OpenCL_C 120
|
||||
%164 = OpString "main"
|
||||
%165 = OpString " __kernel"
|
||||
%167 = OpString "out_data"
|
||||
%170 = OpString "pix_in_block"
|
||||
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
|
||||
OpDecorate %gl_WorkGroupID BuiltIn WorkgroupId
|
||||
OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
|
||||
OpDecorate %_runtimearr_v4uint ArrayStride 16
|
||||
OpMemberDecorate %_struct_23 0 Offset 0
|
||||
OpDecorate %_struct_23 Block
|
||||
OpMemberDecorate %_struct_26 0 Offset 0
|
||||
OpMemberDecorate %_struct_27 0 Offset 0
|
||||
OpDecorate %_struct_27 Block
|
||||
OpDecorate %25 DescriptorSet 0
|
||||
OpDecorate %25 Binding 0
|
||||
OpDecorate %_arr_uint_uint_256 ArrayStride 4
|
||||
OpDecorate %_arr_uchar_uint_1024 ArrayStride 1
|
||||
OpDecorate %15 SpecId 0
|
||||
OpDecorate %16 SpecId 1
|
||||
OpDecorate %17 SpecId 2
|
||||
%uint = OpTypeInt 32 0
|
||||
%uint_256 = OpConstant %uint 256
|
||||
%_arr_uint_uint_256 = OpTypeArray %uint %uint_256
|
||||
%_ptr_Workgroup__arr_uint_uint_256 = OpTypePointer Workgroup %_arr_uint_uint_256
|
||||
%uchar = OpTypeInt 8 0
|
||||
%uint_1024 = OpConstant %uint 1024
|
||||
%_arr_uchar_uint_1024 = OpTypeArray %uchar %uint_1024
|
||||
%_ptr_Workgroup__arr_uchar_uint_1024 = OpTypePointer Workgroup %_arr_uchar_uint_1024
|
||||
%v3uint = OpTypeVector %uint 3
|
||||
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
|
||||
%15 = OpSpecConstant %uint 1
|
||||
%16 = OpSpecConstant %uint 1
|
||||
%17 = OpSpecConstant %uint 1
|
||||
%gl_WorkGroupSize = OpSpecConstantComposite %v3uint %15 %16 %17
|
||||
%_ptr_Private_v3uint = OpTypePointer Private %v3uint
|
||||
%v4uint = OpTypeVector %uint 4
|
||||
%_runtimearr_v4uint = OpTypeRuntimeArray %v4uint
|
||||
%_struct_23 = OpTypeStruct %_runtimearr_v4uint
|
||||
%_ptr_StorageBuffer__struct_23 = OpTypePointer StorageBuffer %_struct_23
|
||||
%_struct_26 = OpTypeStruct %uint
|
||||
%_struct_27 = OpTypeStruct %_struct_26
|
||||
%_ptr_PushConstant__struct_27 = OpTypePointer PushConstant %_struct_27
|
||||
%void = OpTypeVoid
|
||||
%31 = OpTypeFunction %void
|
||||
%_ptr_PushConstant__struct_26 = OpTypePointer PushConstant %_struct_26
|
||||
%uint_0 = OpConstant %uint 0
|
||||
%bool = OpTypeBool
|
||||
%_ptr_Input_uint = OpTypePointer Input %uint
|
||||
%uint_1 = OpConstant %uint 1
|
||||
%uint_255 = OpConstant %uint 255
|
||||
%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
|
||||
%uint_2 = OpConstant %uint 2
|
||||
%uint_10 = OpConstant %uint 10
|
||||
%uint_1020 = OpConstant %uint 1020
|
||||
%v4uchar = OpTypeVector %uchar 4
|
||||
%_ptr_Workgroup_uchar = OpTypePointer Workgroup %uchar
|
||||
%uint_3 = OpConstant %uint 3
|
||||
%uint_264 = OpConstant %uint 264
|
||||
%_ptr_StorageBuffer_v4uint = OpTypePointer StorageBuffer %v4uint
|
||||
%137 = OpUndef %v4uchar
|
||||
%uint_4 = OpConstant %uint 4
|
||||
%5 = OpVariable %_ptr_Workgroup__arr_uint_uint_256 Workgroup
|
||||
%10 = OpVariable %_ptr_Workgroup__arr_uchar_uint_1024 Workgroup
|
||||
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
|
||||
%gl_WorkGroupID = OpVariable %_ptr_Input_v3uint Input
|
||||
%20 = OpVariable %_ptr_Private_v3uint Private %gl_WorkGroupSize
|
||||
%25 = OpVariable %_ptr_StorageBuffer__struct_23 StorageBuffer
|
||||
%29 = OpVariable %_ptr_PushConstant__struct_27 PushConstant
|
||||
%32 = OpFunction %void None %31
|
||||
%33 = OpLabel
|
||||
%36 = OpAccessChain %_ptr_PushConstant__struct_26 %29 %uint_0
|
||||
%37 = OpLoad %_struct_26 %36
|
||||
%38 = OpCompositeExtract %uint %37 0
|
||||
%40 = OpINotEqual %bool %38 %uint_0
|
||||
OpSelectionMerge %105 None
|
||||
OpBranchConditional %40 %43 %105
|
||||
%43 = OpLabel
|
||||
%45 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_0
|
||||
%46 = OpLoad %uint %45
|
||||
%48 = OpAccessChain %_ptr_Input_uint %gl_WorkGroupID %uint_1
|
||||
%49 = OpLoad %uint %48
|
||||
%50 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0
|
||||
%51 = OpLoad %uint %50
|
||||
%52 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1
|
||||
%53 = OpLoad %uint %52
|
||||
%54 = OpIMul %uint %53 %51
|
||||
%55 = OpUDiv %uint %54 %49
|
||||
%57 = OpUMod %uint %55 %uint_255
|
||||
%58 = OpUConvert %uchar %57
|
||||
OpBranch %60
|
||||
%60 = OpLabel
|
||||
%61 = OpPhi %uint %100 %99 %uint_0 %43
|
||||
%62 = OpIMul %uint %61 %38
|
||||
OpLoopMerge %103 %99 None
|
||||
OpBranch %65
|
||||
%65 = OpLabel
|
||||
%66 = OpPhi %uint %93 %65 %uint_0 %60
|
||||
%67 = OpIAdd %uint %66 %62
|
||||
%68 = OpIMul %uint %66 %61
|
||||
%69 = OpIAdd %uint %46 %68
|
||||
%71 = OpAccessChain %_ptr_Workgroup_uint %5 %67
|
||||
OpStore %71 %69
|
||||
%72 = OpIAdd %uint %49 %68
|
||||
%74 = OpShiftLeftLogical %uint %67 %uint_2
|
||||
%76 = OpShiftRightLogical %uint %74 %uint_10
|
||||
%78 = OpBitwiseAnd %uint %74 %uint_1020
|
||||
%80 = OpBitcast %v4uchar %72
|
||||
%81 = OpCompositeExtract %uchar %80 1
|
||||
%82 = OpCompositeExtract %uchar %80 2
|
||||
%83 = OpCompositeExtract %uchar %80 3
|
||||
%85 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %76 %78
|
||||
%86 = OpBitwiseOr %uint %78 %uint_1
|
||||
%87 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %76 %86
|
||||
OpStore %87 %81
|
||||
%88 = OpBitwiseOr %uint %78 %uint_2
|
||||
%89 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %76 %88
|
||||
OpStore %89 %82
|
||||
%91 = OpBitwiseOr %uint %78 %uint_3
|
||||
%92 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %76 %91
|
||||
OpStore %92 %83
|
||||
OpStore %85 %58
|
||||
%93 = OpIAdd %uint %66 %uint_1
|
||||
%94 = OpUGreaterThanEqual %bool %93 %38
|
||||
OpLoopMerge %97 %65 None
|
||||
OpBranchConditional %94 %97 %65
|
||||
%97 = OpLabel
|
||||
OpBranch %99
|
||||
%99 = OpLabel
|
||||
%100 = OpIAdd %uint %61 %uint_1
|
||||
%101 = OpUGreaterThanEqual %bool %100 %38
|
||||
OpBranchConditional %101 %103 %60
|
||||
%103 = OpLabel
|
||||
OpBranch %105
|
||||
%105 = OpLabel
|
||||
OpBranch %107
|
||||
%107 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
OpSelectionMerge %162 None
|
||||
OpBranchConditional %40 %111 %162
|
||||
%111 = OpLabel
|
||||
%112 = OpPhi %uint %157 %156 %uint_0 %107
|
||||
%113 = OpIMul %uint %112 %38
|
||||
OpLoopMerge %160 %156 None
|
||||
OpBranch %116
|
||||
%116 = OpLabel
|
||||
%117 = OpPhi %uint %150 %116 %uint_0 %111
|
||||
%118 = OpIAdd %uint %117 %113
|
||||
%120 = OpAccessChain %_ptr_StorageBuffer_v4uint %25 %uint_0 %118
|
||||
%121 = OpAccessChain %_ptr_Workgroup_uint %5 %118
|
||||
%122 = OpLoad %uint %121
|
||||
%123 = OpShiftLeftLogical %uint %118 %uint_2
|
||||
%124 = OpShiftRightLogical %uint %123 %uint_10
|
||||
%125 = OpBitwiseAnd %uint %123 %uint_1020
|
||||
%126 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %124 %125
|
||||
%127 = OpLoad %uchar %126
|
||||
%128 = OpBitwiseOr %uint %125 %uint_1
|
||||
%129 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %124 %128
|
||||
%130 = OpLoad %uchar %129
|
||||
%131 = OpBitwiseOr %uint %125 %uint_2
|
||||
%132 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %124 %131
|
||||
%133 = OpLoad %uchar %132
|
||||
%134 = OpBitwiseOr %uint %125 %uint_3
|
||||
%135 = OpPtrAccessChain %_ptr_Workgroup_uchar %10 %124 %134
|
||||
%136 = OpLoad %uchar %135
|
||||
%138 = OpCompositeInsert %v4uchar %127 %137 0
|
||||
%139 = OpCompositeInsert %v4uchar %130 %138 1
|
||||
%140 = OpCompositeInsert %v4uchar %133 %139 2
|
||||
%141 = OpCompositeInsert %v4uchar %136 %140 3
|
||||
%142 = OpBitcast %uint %141
|
||||
%143 = OpIAdd %uint %122 %142
|
||||
%144 = OpLoad %v4uint %120
|
||||
%145 = OpCompositeInsert %v4uint %143 %144 0
|
||||
%146 = OpShiftRightLogical %uint %143 %uint_2
|
||||
%147 = OpCompositeInsert %v4uint %146 %145 1
|
||||
%148 = OpShiftRightLogical %uint %143 %uint_3
|
||||
%149 = OpCompositeInsert %v4uint %148 %147 3
|
||||
OpStore %120 %149
|
||||
%150 = OpIAdd %uint %117 %uint_1
|
||||
%151 = OpUGreaterThanEqual %bool %150 %38
|
||||
OpLoopMerge %154 %116 None
|
||||
OpBranchConditional %151 %154 %116
|
||||
%154 = OpLabel
|
||||
OpBranch %156
|
||||
%156 = OpLabel
|
||||
%157 = OpIAdd %uint %112 %uint_1
|
||||
%158 = OpUGreaterThanEqual %bool %157 %38
|
||||
OpBranchConditional %158 %160 %111
|
||||
%160 = OpLabel
|
||||
OpBranch %162
|
||||
%162 = OpLabel
|
||||
OpControlBarrier %uint_2 %uint_2 %uint_264
|
||||
OpReturn
|
||||
OpFunctionEnd
|
||||
%166 = OpExtInst %void %163 Kernel %32 %164 %uint_2 %uint_0 %165
|
||||
%168 = OpExtInst %void %163 ArgumentInfo %167
|
||||
%169 = OpExtInst %void %163 ArgumentStorageBuffer %166 %uint_0 %uint_0 %uint_0 %168
|
||||
%171 = OpExtInst %void %163 ArgumentInfo %170
|
||||
%173 = OpExtInst %void %163 ArgumentPodPushConstant %166 %uint_1 %uint_0 %uint_4 %171
|
||||
%174 = OpExtInst %void %163 SpecConstantWorkgroupSize %uint_0 %uint_1 %uint_2
|
||||
@ -9914,16 +9914,21 @@ void CompilerGLSL::access_chain_internal_append_index(std::string &expr, uint32_
|
||||
if (ptr_chain && access_chain_is_arrayed)
|
||||
{
|
||||
size_t split_pos = expr.find_last_of(']');
|
||||
string expr_front = expr.substr(0, split_pos);
|
||||
string expr_back = expr.substr(split_pos);
|
||||
expr = expr_front + " + " + enclose_expression(idx_expr) + expr_back;
|
||||
}
|
||||
else
|
||||
{
|
||||
expr += "[";
|
||||
expr += idx_expr;
|
||||
expr += "]";
|
||||
size_t enclose_split = expr.find_last_of(')');
|
||||
|
||||
// If we have already enclosed the expression, don't try to be clever, it will break.
|
||||
if (split_pos > enclose_split || enclose_split == string::npos)
|
||||
{
|
||||
string expr_front = expr.substr(0, split_pos);
|
||||
string expr_back = expr.substr(split_pos);
|
||||
expr = expr_front + " + " + enclose_expression(idx_expr) + expr_back;
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
expr += "[";
|
||||
expr += idx_expr;
|
||||
expr += "]";
|
||||
}
|
||||
|
||||
bool CompilerGLSL::access_chain_needs_stage_io_builtin_translation(uint32_t)
|
||||
@ -9958,6 +9963,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
||||
// Start traversing type hierarchy at the proper non-pointer types,
|
||||
// but keep type_id referencing the original pointer for use below.
|
||||
uint32_t type_id = expression_type_id(base);
|
||||
const auto *type = &get_pointee_type(type_id);
|
||||
|
||||
if (!backend.native_pointers)
|
||||
{
|
||||
@ -9967,13 +9973,10 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
||||
// Wrapped buffer reference pointer types will need to poke into the internal "value" member before
|
||||
// continuing the access chain.
|
||||
if (should_dereference(base))
|
||||
{
|
||||
auto &type = get<SPIRType>(type_id);
|
||||
expr = dereference_expression(type, expr);
|
||||
}
|
||||
expr = dereference_expression(get<SPIRType>(type_id), expr);
|
||||
}
|
||||
|
||||
const auto *type = &get_pointee_type(type_id);
|
||||
else if (should_dereference(base) && type->basetype != SPIRType::Struct && !ptr_chain)
|
||||
expr = join("(", dereference_expression(*type, expr), ")");
|
||||
|
||||
bool access_chain_is_arrayed = expr.find_first_of('[') != string::npos;
|
||||
bool row_major_matrix_needs_conversion = is_non_native_row_major_matrix(base);
|
||||
@ -10014,9 +10017,21 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
||||
index &= 0x7fffffffu;
|
||||
}
|
||||
|
||||
// Pointer chains
|
||||
bool ptr_chain_array_entry = ptr_chain && i == 0 && type_is_top_level_array(*type);
|
||||
|
||||
if (ptr_chain_array_entry)
|
||||
{
|
||||
// This is highly unusual code, since normally we'd use plain AccessChain, but it's still allowed.
|
||||
// We are considered to have a pointer to array and one element shifts by one array at a time.
|
||||
// If we use normal array indexing, we'll first decay to pointer, and lose the array-ness,
|
||||
// so we have to take pointer to array explicitly.
|
||||
if (!should_dereference(base))
|
||||
expr = enclose_expression(address_of_expression(expr));
|
||||
}
|
||||
|
||||
if (ptr_chain && i == 0)
|
||||
{
|
||||
// Pointer chains
|
||||
// If we are flattening multidimensional arrays, only create opening bracket on first
|
||||
// array index.
|
||||
if (options.flatten_multidimensional_arrays)
|
||||
@ -10061,6 +10076,12 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
|
||||
}
|
||||
|
||||
access_chain_is_arrayed = true;
|
||||
|
||||
// Explicitly enclose the expression if this is one of the weird pointer-to-array cases.
|
||||
// We don't want any future indexing to add to this array dereference.
|
||||
// Enclosing the expression blocks that and avoids any shenanigans with operand priority.
|
||||
if (ptr_chain_array_entry)
|
||||
expr = join("(", expr, ")");
|
||||
}
|
||||
// Arrays
|
||||
else if (!type->array.empty())
|
||||
|
||||
@ -9863,7 +9863,7 @@ uint32_t CompilerMSL::get_physical_tess_level_array_size(spv::BuiltIn builtin) c
|
||||
bool CompilerMSL::maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs)
|
||||
{
|
||||
// We only care about assignments of an entire array
|
||||
auto &type = expression_type(id_rhs);
|
||||
auto &type = expression_type(id_lhs);
|
||||
if (!type_is_top_level_array(get_pointee_type(type)))
|
||||
return false;
|
||||
|
||||
|
||||
Loading…
Reference in New Issue
Block a user