Welcome to mirror list, hosted at ThFree Co, Russian Federation.

github.com/KhronosGroup/SPIRV-Cross.git - Unnamed repository; edit this file 'description' to name the repository.
summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorHans-Kristian Arntzen <post@arntzen-software.no>2022-01-18 14:39:16 +0300
committerHans-Kristian Arntzen <post@arntzen-software.no>2022-01-18 14:39:16 +0300
commitac46140ba394ba5f85b6cfc72a5f5810084e8191 (patch)
tree2b2253b3486fd10b0126c9d4233a80114c1bdde9
parent48b5a9069f69b2370188a9a9ec4318cc784bb291 (diff)
Test aliased names in declared LUTs.
-rw-r--r--reference/shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp27
-rw-r--r--reference/shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp61
-rw-r--r--reference/shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp16
-rw-r--r--shaders-hlsl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp81
-rw-r--r--shaders-msl-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp81
-rw-r--r--shaders-no-opt/asm/comp/constant-lut-name-aliasing.asm.comp81
-rw-r--r--spirv_hlsl.cpp1
7 files changed, 348 insertions, 0 deletions
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 <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 SSBO
+{
+ int values[1];
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(4u, 4u, 1u);
+
+constant spvUnsafeArray<int, 4> indexable = spvUnsafeArray<int, 4>({ 0, 1, 2, 3 });
+constant spvUnsafeArray<int, 4> indexable_1 = spvUnsafeArray<int, 4>({ 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;