From ac46140ba394ba5f85b6cfc72a5f5810084e8191 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Tue, 18 Jan 2022 12:39:16 +0100 Subject: [PATCH] Test aliased names in declared LUTs. --- .../comp/constant-lut-name-aliasing.asm.comp | 27 +++++++ .../comp/constant-lut-name-aliasing.asm.comp | 61 ++++++++++++++ .../comp/constant-lut-name-aliasing.asm.comp | 16 ++++ .../comp/constant-lut-name-aliasing.asm.comp | 81 +++++++++++++++++++ .../comp/constant-lut-name-aliasing.asm.comp | 81 +++++++++++++++++++ .../comp/constant-lut-name-aliasing.asm.comp | 81 +++++++++++++++++++ spirv_hlsl.cpp | 1 + 7 files changed, 348 insertions(+) create mode 100644 reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp create mode 100644 reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp create mode 100644 reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp create mode 100644 shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp create mode 100644 shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp create mode 100644 shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp diff --git a/reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..d3dc5337 --- /dev/null +++ b/reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -0,0 +1,27 @@ +static const uint3 gl_WorkGroupSize = uint3(4u, 4u, 1u); + +static const int indexable[4] = { 0, 1, 2, 3 }; +static const int indexable_1[4] = { 4, 5, 6, 7 }; + +RWByteAddressBuffer _6 : register(u0); + +static uint3 gl_LocalInvocationID; +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_LocalInvocationID : SV_GroupThreadID; + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + _6.Store(gl_GlobalInvocationID.x * 4 + 0, uint(indexable[gl_LocalInvocationID.x] + indexable_1[gl_LocalInvocationID.y])); +} + +[numthreads(4, 4, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_LocalInvocationID = stage_input.gl_LocalInvocationID; + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} diff --git a/reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..a5b6fc32 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -0,0 +1,61 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" +#pragma clang diagnostic ignored "-Wmissing-braces" + +#include +#include + +using namespace metal; + +template +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 SSBO +{ + int values[1]; +}; + +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 4u, 1u); + +constant spvUnsafeArray indexable = spvUnsafeArray({ 0, 1, 2, 3 }); +constant spvUnsafeArray indexable_1 = spvUnsafeArray({ 4, 5, 6, 7 }); + +kernel void main0(device SSBO& _6 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]]) +{ + _6.values[gl_GlobalInvocationID.x] = indexable[gl_LocalInvocationID.x] + indexable_1[gl_LocalInvocationID.y]; +} + diff --git a/reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..1f43951a --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -0,0 +1,16 @@ +#version 450 +layout(local_size_x = 4, local_size_y = 4, local_size_z = 1) in; + +const int indexable[4] = int[](0, 1, 2, 3); +const int indexable_1[4] = int[](4, 5, 6, 7); + +layout(binding = 0, std430) buffer SSBO +{ + int values[]; +} _6; + +void main() +{ + _6.values[gl_GlobalInvocationID.x] = indexable[gl_LocalInvocationID.x] + indexable_1[gl_LocalInvocationID.y]; +} + diff --git a/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..e1dcb0ef --- /dev/null +++ b/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -0,0 +1,81 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 49 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID + OpExecutionMode %main LocalSize 4 4 1 + OpSource GLSL 450 + OpName %main "main" + OpName %SSBO "SSBO" + OpMemberName %SSBO 0 "values" + OpName %_ "" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %gl_LocalInvocationID "gl_LocalInvocationID" + OpName %indexable "indexable" + OpName %indexable_0 "indexable" + OpName %25 "indexable" + OpName %38 "indexable" + OpDecorate %_runtimearr_int ArrayStride 4 + OpMemberDecorate %SSBO 0 Offset 0 + OpDecorate %SSBO BufferBlock + OpDecorate %_ DescriptorSet 0 + OpDecorate %_ Binding 0 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 +%_runtimearr_int = OpTypeRuntimeArray %int + %SSBO = OpTypeStruct %_runtimearr_int +%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO + %_ = OpVariable %_ptr_Uniform_SSBO Uniform + %int_0 = OpConstant %int 0 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %uint_0 = OpConstant %uint 0 +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_4 = OpConstant %uint 4 +%_arr_int_uint_4 = OpTypeArray %int %uint_4 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %25 = OpConstantComposite %_arr_int_uint_4 %int_0 %int_1 %int_2 %int_3 +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input +%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 +%_ptr_Function_int = OpTypePointer Function %int + %int_4 = OpConstant %int 4 + %int_5 = OpConstant %int 5 + %int_6 = OpConstant %int 6 + %int_7 = OpConstant %int 7 + %38 = OpConstantComposite %_arr_int_uint_4 %int_4 %int_5 %int_6 %int_7 + %uint_1 = OpConstant %uint 1 +%_ptr_Uniform_int = OpTypePointer Uniform %int +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %indexable = OpVariable %_ptr_Function__arr_int_uint_4 Function +%indexable_0 = OpVariable %_ptr_Function__arr_int_uint_4 Function + %18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %19 = OpLoad %uint %18 + %27 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %28 = OpLoad %uint %27 + OpStore %indexable %25 + %32 = OpAccessChain %_ptr_Function_int %indexable %28 + %33 = OpLoad %int %32 + %40 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1 + %41 = OpLoad %uint %40 + OpStore %indexable_0 %38 + %43 = OpAccessChain %_ptr_Function_int %indexable_0 %41 + %44 = OpLoad %int %43 + %45 = OpIAdd %int %33 %44 + %47 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19 + OpStore %47 %45 + OpReturn + OpFunctionEnd diff --git a/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..e1dcb0ef --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -0,0 +1,81 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 49 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID + OpExecutionMode %main LocalSize 4 4 1 + OpSource GLSL 450 + OpName %main "main" + OpName %SSBO "SSBO" + OpMemberName %SSBO 0 "values" + OpName %_ "" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %gl_LocalInvocationID "gl_LocalInvocationID" + OpName %indexable "indexable" + OpName %indexable_0 "indexable" + OpName %25 "indexable" + OpName %38 "indexable" + OpDecorate %_runtimearr_int ArrayStride 4 + OpMemberDecorate %SSBO 0 Offset 0 + OpDecorate %SSBO BufferBlock + OpDecorate %_ DescriptorSet 0 + OpDecorate %_ Binding 0 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 +%_runtimearr_int = OpTypeRuntimeArray %int + %SSBO = OpTypeStruct %_runtimearr_int +%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO + %_ = OpVariable %_ptr_Uniform_SSBO Uniform + %int_0 = OpConstant %int 0 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %uint_0 = OpConstant %uint 0 +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_4 = OpConstant %uint 4 +%_arr_int_uint_4 = OpTypeArray %int %uint_4 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %25 = OpConstantComposite %_arr_int_uint_4 %int_0 %int_1 %int_2 %int_3 +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input +%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 +%_ptr_Function_int = OpTypePointer Function %int + %int_4 = OpConstant %int 4 + %int_5 = OpConstant %int 5 + %int_6 = OpConstant %int 6 + %int_7 = OpConstant %int 7 + %38 = OpConstantComposite %_arr_int_uint_4 %int_4 %int_5 %int_6 %int_7 + %uint_1 = OpConstant %uint 1 +%_ptr_Uniform_int = OpTypePointer Uniform %int +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %indexable = OpVariable %_ptr_Function__arr_int_uint_4 Function +%indexable_0 = OpVariable %_ptr_Function__arr_int_uint_4 Function + %18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %19 = OpLoad %uint %18 + %27 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %28 = OpLoad %uint %27 + OpStore %indexable %25 + %32 = OpAccessChain %_ptr_Function_int %indexable %28 + %33 = OpLoad %int %32 + %40 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1 + %41 = OpLoad %uint %40 + OpStore %indexable_0 %38 + %43 = OpAccessChain %_ptr_Function_int %indexable_0 %41 + %44 = OpLoad %int %43 + %45 = OpIAdd %int %33 %44 + %47 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19 + OpStore %47 %45 + OpReturn + OpFunctionEnd diff --git a/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp b/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp new file mode 100644 index 00000000..e1dcb0ef --- /dev/null +++ b/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp @@ -0,0 +1,81 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 49 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID + OpExecutionMode %main LocalSize 4 4 1 + OpSource GLSL 450 + OpName %main "main" + OpName %SSBO "SSBO" + OpMemberName %SSBO 0 "values" + OpName %_ "" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %gl_LocalInvocationID "gl_LocalInvocationID" + OpName %indexable "indexable" + OpName %indexable_0 "indexable" + OpName %25 "indexable" + OpName %38 "indexable" + OpDecorate %_runtimearr_int ArrayStride 4 + OpMemberDecorate %SSBO 0 Offset 0 + OpDecorate %SSBO BufferBlock + OpDecorate %_ DescriptorSet 0 + OpDecorate %_ Binding 0 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId + OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 +%_runtimearr_int = OpTypeRuntimeArray %int + %SSBO = OpTypeStruct %_runtimearr_int +%_ptr_Uniform_SSBO = OpTypePointer Uniform %SSBO + %_ = OpVariable %_ptr_Uniform_SSBO Uniform + %int_0 = OpConstant %int 0 + %uint = OpTypeInt 32 0 + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input + %uint_0 = OpConstant %uint 0 +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_4 = OpConstant %uint 4 +%_arr_int_uint_4 = OpTypeArray %int %uint_4 + %int_1 = OpConstant %int 1 + %int_2 = OpConstant %int 2 + %int_3 = OpConstant %int 3 + %25 = OpConstantComposite %_arr_int_uint_4 %int_0 %int_1 %int_2 %int_3 +%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input +%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4 +%_ptr_Function_int = OpTypePointer Function %int + %int_4 = OpConstant %int 4 + %int_5 = OpConstant %int 5 + %int_6 = OpConstant %int 6 + %int_7 = OpConstant %int 7 + %38 = OpConstantComposite %_arr_int_uint_4 %int_4 %int_5 %int_6 %int_7 + %uint_1 = OpConstant %uint 1 +%_ptr_Uniform_int = OpTypePointer Uniform %int +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_4 %uint_4 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %indexable = OpVariable %_ptr_Function__arr_int_uint_4 Function +%indexable_0 = OpVariable %_ptr_Function__arr_int_uint_4 Function + %18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %19 = OpLoad %uint %18 + %27 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_0 + %28 = OpLoad %uint %27 + OpStore %indexable %25 + %32 = OpAccessChain %_ptr_Function_int %indexable %28 + %33 = OpLoad %int %32 + %40 = OpAccessChain %_ptr_Input_uint %gl_LocalInvocationID %uint_1 + %41 = OpLoad %uint %40 + OpStore %indexable_0 %38 + %43 = OpAccessChain %_ptr_Function_int %indexable_0 %41 + %44 = OpLoad %int %43 + %45 = OpIAdd %int %33 %44 + %47 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19 + OpStore %47 %45 + OpReturn + OpFunctionEnd diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp index c09db387..3d834774 100644 --- a/spirv_hlsl.cpp +++ b/spirv_hlsl.cpp @@ -1167,6 +1167,7 @@ void CompilerHLSL::emit_composite_constants() if (type.basetype == SPIRType::Struct || !type.array.empty()) { + add_resource_name(c.self); auto name = to_name(c.self); statement("static const ", variable_decl(type, name), " = ", constant_expression(c), ";"); emitted = true;