diff options
16 files changed, 748 insertions, 1 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-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp b/reference/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp new file mode 100644 index 00000000..423beee6 --- /dev/null +++ b/reference/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp @@ -0,0 +1,51 @@ +#ifndef SPIRV_CROSS_CONSTANT_ID_0 +#define SPIRV_CROSS_CONSTANT_ID_0 0 +#endif +static const int A = SPIRV_CROSS_CONSTANT_ID_0; +#ifndef SPIRV_CROSS_CONSTANT_ID_1 +#define SPIRV_CROSS_CONSTANT_ID_1 1 +#endif +static const int A_1 = SPIRV_CROSS_CONSTANT_ID_1; +#ifndef SPIRV_CROSS_CONSTANT_ID_2 +#define SPIRV_CROSS_CONSTANT_ID_2 2 +#endif +static const int A_2 = SPIRV_CROSS_CONSTANT_ID_2; +#ifndef SPIRV_CROSS_CONSTANT_ID_3 +#define SPIRV_CROSS_CONSTANT_ID_3 3 +#endif +static const int A_3 = SPIRV_CROSS_CONSTANT_ID_3; +#ifndef SPIRV_CROSS_CONSTANT_ID_4 +#define SPIRV_CROSS_CONSTANT_ID_4 4 +#endif +static const int A_4 = SPIRV_CROSS_CONSTANT_ID_4; +#ifndef SPIRV_CROSS_CONSTANT_ID_5 +#define SPIRV_CROSS_CONSTANT_ID_5 5 +#endif +static const int A_5 = SPIRV_CROSS_CONSTANT_ID_5; +static const int A_6 = (A - A_1); +static const int A_7 = (A_6 - A_2); +static const int A_8 = (A_7 - A_3); +static const int A_9 = (A_8 - A_4); +static const int A_10 = (A_9 - A_5); +static const int A_11 = (A_10 + A_5); +static const uint3 gl_WorkGroupSize = uint3(1u, 1u, 1u); + +RWByteAddressBuffer _5 : register(u0); + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + _5.Store(gl_GlobalInvocationID.x * 4 + 0, uint(A_11)); +} + +[numthreads(1, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + 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-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp b/reference/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp new file mode 100644 index 00000000..dda85050 --- /dev/null +++ b/reference/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp @@ -0,0 +1,35 @@ +#include <metal_stdlib> +#include <simd/simd.h> + +using namespace metal; + +struct SSBO +{ + int values[1]; +}; + +constant int A_tmp [[function_constant(0)]]; +constant int A = is_function_constant_defined(A_tmp) ? A_tmp : 0; +constant int A_1_tmp [[function_constant(1)]]; +constant int A_1 = is_function_constant_defined(A_1_tmp) ? A_1_tmp : 1; +constant int A_2_tmp [[function_constant(2)]]; +constant int A_2 = is_function_constant_defined(A_2_tmp) ? A_2_tmp : 2; +constant int A_3_tmp [[function_constant(3)]]; +constant int A_3 = is_function_constant_defined(A_3_tmp) ? A_3_tmp : 3; +constant int A_4_tmp [[function_constant(4)]]; +constant int A_4 = is_function_constant_defined(A_4_tmp) ? A_4_tmp : 4; +constant int A_5_tmp [[function_constant(5)]]; +constant int A_5 = is_function_constant_defined(A_5_tmp) ? A_5_tmp : 5; +constant int A_6 = (A - A_1); +constant int A_7 = (A_6 - A_2); +constant int A_8 = (A_7 - A_3); +constant int A_9 = (A_8 - A_4); +constant int A_10 = (A_9 - A_5); +constant int A_11 = (A_10 + A_5); +constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u); + +kernel void main0(device SSBO& _5 [[buffer(0)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + _5.values[gl_GlobalInvocationID.x] = A_11; +} + 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/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp new file mode 100644 index 00000000..e7b9dbf3 --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp @@ -0,0 +1,44 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +#ifndef SPIRV_CROSS_CONSTANT_ID_0 +#define SPIRV_CROSS_CONSTANT_ID_0 0 +#endif +const int A = SPIRV_CROSS_CONSTANT_ID_0; +#ifndef SPIRV_CROSS_CONSTANT_ID_1 +#define SPIRV_CROSS_CONSTANT_ID_1 1 +#endif +const int A_1 = SPIRV_CROSS_CONSTANT_ID_1; +#ifndef SPIRV_CROSS_CONSTANT_ID_2 +#define SPIRV_CROSS_CONSTANT_ID_2 2 +#endif +const int A_2 = SPIRV_CROSS_CONSTANT_ID_2; +#ifndef SPIRV_CROSS_CONSTANT_ID_3 +#define SPIRV_CROSS_CONSTANT_ID_3 3 +#endif +const int A_3 = SPIRV_CROSS_CONSTANT_ID_3; +#ifndef SPIRV_CROSS_CONSTANT_ID_4 +#define SPIRV_CROSS_CONSTANT_ID_4 4 +#endif +const int A_4 = SPIRV_CROSS_CONSTANT_ID_4; +#ifndef SPIRV_CROSS_CONSTANT_ID_5 +#define SPIRV_CROSS_CONSTANT_ID_5 5 +#endif +const int A_5 = SPIRV_CROSS_CONSTANT_ID_5; +const int A_6 = (A - A_1); +const int A_7 = (A_6 - A_2); +const int A_8 = (A_7 - A_3); +const int A_9 = (A_8 - A_4); +const int A_10 = (A_9 - A_5); +const int A_11 = (A_10 + A_5); + +layout(binding = 0, std430) buffer SSBO +{ + int values[]; +} _5; + +void main() +{ + _5.values[gl_GlobalInvocationID.x] = A_11; +} + diff --git a/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp.vk b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp.vk new file mode 100644 index 00000000..c31d0787 --- /dev/null +++ b/reference/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp.vk @@ -0,0 +1,26 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(constant_id = 0) const int A = 0; +layout(constant_id = 1) const int A_1 = 1; +layout(constant_id = 2) const int A_2 = 2; +layout(constant_id = 3) const int A_3 = 3; +layout(constant_id = 4) const int A_4 = 4; +layout(constant_id = 5) const int A_5 = 5; +const int A_6 = (A - A_1); +const int A_7 = (A_6 - A_2); +const int A_8 = (A_7 - A_3); +const int A_9 = (A_8 - A_4); +const int A_10 = (A_9 - A_5); +const int A_11 = (A_10 + A_5); + +layout(set = 0, binding = 0, std430) buffer SSBO +{ + int values[]; +} _5; + +void main() +{ + _5.values[gl_GlobalInvocationID.x] = A_11; +} + 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-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp b/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp new file mode 100644 index 00000000..b4e622ba --- /dev/null +++ b/shaders-hlsl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp @@ -0,0 +1,78 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 35 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %main "main" + OpName %SSBO "SSBO" + OpMemberName %SSBO 0 "values" + OpName %_ "" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %A "A" + OpName %B "A" + OpName %C "A" + OpName %D "A" + OpName %E "A" + OpName %F "A" + OpName %G "A" + OpName %H "A" + OpName %I "A" + OpName %J "A" + OpName %K "A" + OpName %L "A" + 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 %A SpecId 0 + OpDecorate %B SpecId 1 + OpDecorate %C SpecId 2 + OpDecorate %D SpecId 3 + OpDecorate %E SpecId 4 + OpDecorate %F SpecId 5 + 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 + %A = OpSpecConstant %int 0 + %B = OpSpecConstant %int 1 + %C = OpSpecConstant %int 2 + %D = OpSpecConstant %int 3 + %E = OpSpecConstant %int 4 + %F = OpSpecConstant %int 5 + %G = OpSpecConstantOp %int ISub %A %B + %H = OpSpecConstantOp %int ISub %G %C + %I = OpSpecConstantOp %int ISub %H %D + %J = OpSpecConstantOp %int ISub %I %E + %K = OpSpecConstantOp %int ISub %J %F + %L = OpSpecConstantOp %int IAdd %K %F +%_ptr_Uniform_int = OpTypePointer Uniform %int + %uint_1 = OpConstant %uint 1 +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %19 = OpLoad %uint %18 + %32 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19 + OpStore %32 %L + 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-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp b/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp new file mode 100644 index 00000000..b4e622ba --- /dev/null +++ b/shaders-msl-no-opt/asm/comp/spec-constant-name-aliasing.asm.comp @@ -0,0 +1,78 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 35 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %main "main" + OpName %SSBO "SSBO" + OpMemberName %SSBO 0 "values" + OpName %_ "" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %A "A" + OpName %B "A" + OpName %C "A" + OpName %D "A" + OpName %E "A" + OpName %F "A" + OpName %G "A" + OpName %H "A" + OpName %I "A" + OpName %J "A" + OpName %K "A" + OpName %L "A" + 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 %A SpecId 0 + OpDecorate %B SpecId 1 + OpDecorate %C SpecId 2 + OpDecorate %D SpecId 3 + OpDecorate %E SpecId 4 + OpDecorate %F SpecId 5 + 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 + %A = OpSpecConstant %int 0 + %B = OpSpecConstant %int 1 + %C = OpSpecConstant %int 2 + %D = OpSpecConstant %int 3 + %E = OpSpecConstant %int 4 + %F = OpSpecConstant %int 5 + %G = OpSpecConstantOp %int ISub %A %B + %H = OpSpecConstantOp %int ISub %G %C + %I = OpSpecConstantOp %int ISub %H %D + %J = OpSpecConstantOp %int ISub %I %E + %K = OpSpecConstantOp %int ISub %J %F + %L = OpSpecConstantOp %int IAdd %K %F +%_ptr_Uniform_int = OpTypePointer Uniform %int + %uint_1 = OpConstant %uint 1 +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %19 = OpLoad %uint %18 + %32 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19 + OpStore %32 %L + 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/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp b/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp new file mode 100644 index 00000000..b4e622ba --- /dev/null +++ b/shaders-no-opt/asm/comp/spec-constant-name-aliasing.vk.asm.comp @@ -0,0 +1,78 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 10 +; Bound: 35 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID + OpExecutionMode %main LocalSize 1 1 1 + OpSource GLSL 450 + OpName %main "main" + OpName %SSBO "SSBO" + OpMemberName %SSBO 0 "values" + OpName %_ "" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %A "A" + OpName %B "A" + OpName %C "A" + OpName %D "A" + OpName %E "A" + OpName %F "A" + OpName %G "A" + OpName %H "A" + OpName %I "A" + OpName %J "A" + OpName %K "A" + OpName %L "A" + 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 %A SpecId 0 + OpDecorate %B SpecId 1 + OpDecorate %C SpecId 2 + OpDecorate %D SpecId 3 + OpDecorate %E SpecId 4 + OpDecorate %F SpecId 5 + 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 + %A = OpSpecConstant %int 0 + %B = OpSpecConstant %int 1 + %C = OpSpecConstant %int 2 + %D = OpSpecConstant %int 3 + %E = OpSpecConstant %int 4 + %F = OpSpecConstant %int 5 + %G = OpSpecConstantOp %int ISub %A %B + %H = OpSpecConstantOp %int ISub %G %C + %I = OpSpecConstantOp %int ISub %H %D + %J = OpSpecConstantOp %int ISub %I %E + %K = OpSpecConstantOp %int ISub %J %F + %L = OpSpecConstantOp %int IAdd %K %F +%_ptr_Uniform_int = OpTypePointer Uniform %int + %uint_1 = OpConstant %uint 1 +%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1 + %main = OpFunction %void None %3 + %5 = OpLabel + %18 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %19 = OpLoad %uint %18 + %32 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %19 + OpStore %32 %L + OpReturn + OpFunctionEnd diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index acbeb59a..3cdb742f 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -2686,6 +2686,7 @@ string CompilerGLSL::constant_value_macro_name(uint32_t id) void CompilerGLSL::emit_specialization_constant_op(const SPIRConstantOp &constant) { auto &type = get<SPIRType>(constant.basetype); + add_resource_name(constant.self); auto name = to_name(constant.self); statement("const ", variable_decl(type, name), " = ", constant_op_expression(constant), ";"); } @@ -2713,7 +2714,6 @@ int CompilerGLSL::get_constant_mapping_to_workgroup_component(const SPIRConstant void CompilerGLSL::emit_constant(const SPIRConstant &constant) { auto &type = get<SPIRType>(constant.constant_type); - auto name = to_name(constant.self); SpecializationConstant wg_x, wg_y, wg_z; ID workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z); @@ -2740,6 +2740,9 @@ void CompilerGLSL::emit_constant(const SPIRConstant &constant) return; } + add_resource_name(constant.self); + auto name = to_name(constant.self); + // Only scalars have constant IDs. if (has_decoration(constant.self, DecorationSpecId)) { diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp index efd8f99f..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; @@ -1213,6 +1214,7 @@ void CompilerHLSL::emit_specialization_constants_and_structs() else if (c.specialization) { auto &type = get<SPIRType>(c.constant_type); + add_resource_name(c.self); auto name = to_name(c.self); if (has_decoration(c.self, DecorationSpecId)) @@ -1236,6 +1238,7 @@ void CompilerHLSL::emit_specialization_constants_and_structs() { auto &c = id.get<SPIRConstantOp>(); auto &type = get<SPIRType>(c.basetype); + add_resource_name(c.self); auto name = to_name(c.self); statement("static const ", variable_decl(type, name), " = ", constant_op_expression(c), ";"); emitted = true; diff --git a/spirv_msl.cpp b/spirv_msl.cpp index e8af45f3..714a97f5 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -6565,6 +6565,7 @@ void CompilerMSL::declare_constant_arrays() // link into Metal libraries. This is hacky. if (!type.array.empty() && (!fully_inlined || is_scalar(type) || is_vector(type))) { + add_resource_name(c.self); auto name = to_name(c.self); statement(inject_top_level_storage_qualifier(variable_decl(type, name), "constant"), " = ", constant_expression(c), ";"); @@ -6596,6 +6597,7 @@ void CompilerMSL::declare_complex_constant_arrays() auto &type = this->get<SPIRType>(c.constant_type); if (!type.array.empty() && !(is_scalar(type) || is_vector(type))) { + add_resource_name(c.self); auto name = to_name(c.self); statement("", variable_decl(type, name), " = ", constant_expression(c), ";"); emitted = true; @@ -6674,6 +6676,7 @@ void CompilerMSL::emit_specialization_constants_and_structs() { auto &type = get<SPIRType>(c.constant_type); string sc_type_name = type_to_glsl(type); + add_resource_name(c.self); string sc_name = to_name(c.self); string sc_tmp_name = sc_name + "_tmp"; @@ -6716,6 +6719,7 @@ void CompilerMSL::emit_specialization_constants_and_structs() { auto &c = id.get<SPIRConstantOp>(); auto &type = get<SPIRType>(c.basetype); + add_resource_name(c.self); auto name = to_name(c.self); statement("constant ", variable_decl(type, name), " = ", constant_op_expression(c), ";"); emitted = true; |