diff --git a/reference/opt/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp b/reference/opt/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..a12274c0 --- /dev/null +++ b/reference/opt/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,45 @@ +struct A +{ + int a; + int b; +}; + +struct A_1 +{ + int a; + int b; +}; + +RWByteAddressBuffer C1 : register(u1); +cbuffer C2 : register(b2) +{ + A_1 C2_1_Data[1024] : packoffset(c0); +}; + +RWByteAddressBuffer C3 : register(u0); +cbuffer B : register(b3) +{ + A_1 C4_Data[1024] : packoffset(c0); +}; + + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + C1.Store(gl_GlobalInvocationID.x * 8 + 0, uint(C2_1_Data[gl_GlobalInvocationID.x].a)); + C1.Store(gl_GlobalInvocationID.x * 8 + 4, uint(C2_1_Data[gl_GlobalInvocationID.x].b)); + C3.Store(gl_GlobalInvocationID.x * 8 + 0, uint(C4_Data[gl_GlobalInvocationID.x].a)); + C3.Store(gl_GlobalInvocationID.x * 8 + 4, uint(C4_Data[gl_GlobalInvocationID.x].b)); +} + +[numthreads(1, 1, 1)] +void main(SPIRV_Cross_Input stage_input) +{ + gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID; + comp_main(); +} diff --git a/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp b/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..95f2717b --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,45 @@ +#include +#include + +using namespace metal; + +struct A +{ + int a; + int b; +}; + +struct A_1 +{ + A Data[1]; +}; + +struct A_2 +{ + int a; + int b; +}; + +struct A_3 +{ + A_2 Data[1024]; +}; + +struct B +{ + A Data[1]; +}; + +struct B_1 +{ + A_2 Data[1024]; +}; + +kernel void main0(device B& C3 [[buffer(0)]], device A_1& C1 [[buffer(1)]], constant A_3& C2 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a; + C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b; + C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a; + C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b; +} + diff --git a/reference/opt/shaders/asm/comp/block-name-alias-global.asm.comp b/reference/opt/shaders/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..08fccbcd --- /dev/null +++ b/reference/opt/shaders/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,43 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct A +{ + int a; + int b; +}; + +struct A_1 +{ + int a; + int b; +}; + +layout(binding = 1, std430) buffer C1 +{ + A Data[]; +} C1_1; + +layout(binding = 2, std140) uniform C2 +{ + A_1 Data[1024]; +} C2_1; + +layout(binding = 0, std430) buffer B +{ + A Data[]; +} C3; + +layout(binding = 3, std140) uniform B +{ + A_1 Data[1024]; +} C4; + +void main() +{ + C1_1.Data[gl_GlobalInvocationID.x].a = C2_1.Data[gl_GlobalInvocationID.x].a; + C1_1.Data[gl_GlobalInvocationID.x].b = C2_1.Data[gl_GlobalInvocationID.x].b; + C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a; + C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b; +} + diff --git a/reference/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp b/reference/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..a12274c0 --- /dev/null +++ b/reference/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,45 @@ +struct A +{ + int a; + int b; +}; + +struct A_1 +{ + int a; + int b; +}; + +RWByteAddressBuffer C1 : register(u1); +cbuffer C2 : register(b2) +{ + A_1 C2_1_Data[1024] : packoffset(c0); +}; + +RWByteAddressBuffer C3 : register(u0); +cbuffer B : register(b3) +{ + A_1 C4_Data[1024] : packoffset(c0); +}; + + +static uint3 gl_GlobalInvocationID; +struct SPIRV_Cross_Input +{ + uint3 gl_GlobalInvocationID : SV_DispatchThreadID; +}; + +void comp_main() +{ + C1.Store(gl_GlobalInvocationID.x * 8 + 0, uint(C2_1_Data[gl_GlobalInvocationID.x].a)); + C1.Store(gl_GlobalInvocationID.x * 8 + 4, uint(C2_1_Data[gl_GlobalInvocationID.x].b)); + C3.Store(gl_GlobalInvocationID.x * 8 + 0, uint(C4_Data[gl_GlobalInvocationID.x].a)); + C3.Store(gl_GlobalInvocationID.x * 8 + 4, uint(C4_Data[gl_GlobalInvocationID.x].b)); +} + +[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/asm/comp/block-name-alias-global.asm.comp b/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..95f2717b --- /dev/null +++ b/reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,45 @@ +#include +#include + +using namespace metal; + +struct A +{ + int a; + int b; +}; + +struct A_1 +{ + A Data[1]; +}; + +struct A_2 +{ + int a; + int b; +}; + +struct A_3 +{ + A_2 Data[1024]; +}; + +struct B +{ + A Data[1]; +}; + +struct B_1 +{ + A_2 Data[1024]; +}; + +kernel void main0(device B& C3 [[buffer(0)]], device A_1& C1 [[buffer(1)]], constant A_3& C2 [[buffer(2)]], constant B_1& C4 [[buffer(3)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + C1.Data[gl_GlobalInvocationID.x].a = C2.Data[gl_GlobalInvocationID.x].a; + C1.Data[gl_GlobalInvocationID.x].b = C2.Data[gl_GlobalInvocationID.x].b; + C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a; + C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b; +} + diff --git a/reference/shaders/asm/comp/block-name-alias-global.asm.comp b/reference/shaders/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..08fccbcd --- /dev/null +++ b/reference/shaders/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,43 @@ +#version 450 +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +struct A +{ + int a; + int b; +}; + +struct A_1 +{ + int a; + int b; +}; + +layout(binding = 1, std430) buffer C1 +{ + A Data[]; +} C1_1; + +layout(binding = 2, std140) uniform C2 +{ + A_1 Data[1024]; +} C2_1; + +layout(binding = 0, std430) buffer B +{ + A Data[]; +} C3; + +layout(binding = 3, std140) uniform B +{ + A_1 Data[1024]; +} C4; + +void main() +{ + C1_1.Data[gl_GlobalInvocationID.x].a = C2_1.Data[gl_GlobalInvocationID.x].a; + C1_1.Data[gl_GlobalInvocationID.x].b = C2_1.Data[gl_GlobalInvocationID.x].b; + C3.Data[gl_GlobalInvocationID.x].a = C4.Data[gl_GlobalInvocationID.x].a; + C3.Data[gl_GlobalInvocationID.x].b = C4.Data[gl_GlobalInvocationID.x].b; +} + diff --git a/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp b/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..85f6cc04 --- /dev/null +++ b/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,119 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 59 +; 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 %Foo "A" + OpMemberName %Foo 0 "a" + OpMemberName %Foo 1 "b" + OpName %A "A" + OpMemberName %A 0 "Data" + OpName %C1 "C1" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %Foo_0 "A" + OpMemberName %Foo_0 0 "a" + OpMemberName %Foo_0 1 "b" + OpName %A_0 "A" + OpMemberName %A_0 0 "Data" + OpName %C2 "C2" + OpName %B "B" + OpMemberName %B 0 "Data" + OpName %C3 "C3" + OpName %B_0 "B" + OpMemberName %B_0 0 "Data" + OpName %C4 "C4" + OpMemberDecorate %Foo 0 Offset 0 + OpMemberDecorate %Foo 1 Offset 4 + OpDecorate %_runtimearr_Foo ArrayStride 8 + OpMemberDecorate %A 0 Offset 0 + OpDecorate %A BufferBlock + OpDecorate %C1 DescriptorSet 0 + OpDecorate %C1 Binding 1 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpMemberDecorate %Foo_0 0 Offset 0 + OpMemberDecorate %Foo_0 1 Offset 4 + OpDecorate %_arr_Foo_0_uint_1024 ArrayStride 16 + OpMemberDecorate %A_0 0 Offset 0 + OpDecorate %A_0 Block + OpDecorate %C2 DescriptorSet 0 + OpDecorate %C2 Binding 2 + OpDecorate %_runtimearr_Foo_0 ArrayStride 8 + OpMemberDecorate %B 0 Offset 0 + OpDecorate %B BufferBlock + OpDecorate %C3 DescriptorSet 0 + OpDecorate %C3 Binding 0 + OpDecorate %_arr_Foo_0_uint_1024_0 ArrayStride 16 + OpMemberDecorate %B_0 0 Offset 0 + OpDecorate %B_0 Block + OpDecorate %C4 DescriptorSet 0 + OpDecorate %C4 Binding 3 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %Foo = OpTypeStruct %int %int +%_runtimearr_Foo = OpTypeRuntimeArray %Foo + %A = OpTypeStruct %_runtimearr_Foo +%_ptr_Uniform_A = OpTypePointer Uniform %A + %C1 = OpVariable %_ptr_Uniform_A 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 + %Foo_0 = OpTypeStruct %int %int + %uint_1024 = OpConstant %uint 1024 +%_arr_Foo_0_uint_1024 = OpTypeArray %Foo_0 %uint_1024 + %A_0 = OpTypeStruct %_arr_Foo_0_uint_1024 +%_ptr_Uniform_A_0 = OpTypePointer Uniform %A_0 + %C2 = OpVariable %_ptr_Uniform_A_0 Uniform +%_ptr_Uniform_Foo_0 = OpTypePointer Uniform %Foo_0 +%_ptr_Uniform_Foo = OpTypePointer Uniform %Foo +%_ptr_Uniform_int = OpTypePointer Uniform %int + %int_1 = OpConstant %int 1 +%_runtimearr_Foo_0 = OpTypeRuntimeArray %Foo + %B = OpTypeStruct %_runtimearr_Foo_0 +%_ptr_Uniform_B = OpTypePointer Uniform %B + %C3 = OpVariable %_ptr_Uniform_B Uniform +%_arr_Foo_0_uint_1024_0 = OpTypeArray %Foo_0 %uint_1024 + %B_0 = OpTypeStruct %_arr_Foo_0_uint_1024_0 +%_ptr_Uniform_B_0 = OpTypePointer Uniform %B_0 + %C4 = OpVariable %_ptr_Uniform_B_0 Uniform + %main = OpFunction %void None %3 + %5 = OpLabel + %19 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %20 = OpLoad %uint %19 + %27 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %28 = OpLoad %uint %27 + %30 = OpAccessChain %_ptr_Uniform_Foo_0 %C2 %int_0 %28 + %31 = OpLoad %Foo_0 %30 + %33 = OpAccessChain %_ptr_Uniform_Foo %C1 %int_0 %20 + %34 = OpCompositeExtract %int %31 0 + %36 = OpAccessChain %_ptr_Uniform_int %33 %int_0 + OpStore %36 %34 + %37 = OpCompositeExtract %int %31 1 + %39 = OpAccessChain %_ptr_Uniform_int %33 %int_1 + OpStore %39 %37 + %44 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %45 = OpLoad %uint %44 + %50 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %51 = OpLoad %uint %50 + %52 = OpAccessChain %_ptr_Uniform_Foo_0 %C4 %int_0 %51 + %53 = OpLoad %Foo_0 %52 + %54 = OpAccessChain %_ptr_Uniform_Foo %C3 %int_0 %45 + %55 = OpCompositeExtract %int %53 0 + %56 = OpAccessChain %_ptr_Uniform_int %54 %int_0 + OpStore %56 %55 + %57 = OpCompositeExtract %int %53 1 + %58 = OpAccessChain %_ptr_Uniform_int %54 %int_1 + OpStore %58 %57 + OpReturn + OpFunctionEnd diff --git a/shaders-msl/asm/comp/block-name-alias-global.asm.comp b/shaders-msl/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..85f6cc04 --- /dev/null +++ b/shaders-msl/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,119 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 59 +; 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 %Foo "A" + OpMemberName %Foo 0 "a" + OpMemberName %Foo 1 "b" + OpName %A "A" + OpMemberName %A 0 "Data" + OpName %C1 "C1" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %Foo_0 "A" + OpMemberName %Foo_0 0 "a" + OpMemberName %Foo_0 1 "b" + OpName %A_0 "A" + OpMemberName %A_0 0 "Data" + OpName %C2 "C2" + OpName %B "B" + OpMemberName %B 0 "Data" + OpName %C3 "C3" + OpName %B_0 "B" + OpMemberName %B_0 0 "Data" + OpName %C4 "C4" + OpMemberDecorate %Foo 0 Offset 0 + OpMemberDecorate %Foo 1 Offset 4 + OpDecorate %_runtimearr_Foo ArrayStride 8 + OpMemberDecorate %A 0 Offset 0 + OpDecorate %A BufferBlock + OpDecorate %C1 DescriptorSet 0 + OpDecorate %C1 Binding 1 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpMemberDecorate %Foo_0 0 Offset 0 + OpMemberDecorate %Foo_0 1 Offset 4 + OpDecorate %_arr_Foo_0_uint_1024 ArrayStride 16 + OpMemberDecorate %A_0 0 Offset 0 + OpDecorate %A_0 Block + OpDecorate %C2 DescriptorSet 0 + OpDecorate %C2 Binding 2 + OpDecorate %_runtimearr_Foo_0 ArrayStride 8 + OpMemberDecorate %B 0 Offset 0 + OpDecorate %B BufferBlock + OpDecorate %C3 DescriptorSet 0 + OpDecorate %C3 Binding 0 + OpDecorate %_arr_Foo_0_uint_1024_0 ArrayStride 16 + OpMemberDecorate %B_0 0 Offset 0 + OpDecorate %B_0 Block + OpDecorate %C4 DescriptorSet 0 + OpDecorate %C4 Binding 3 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %Foo = OpTypeStruct %int %int +%_runtimearr_Foo = OpTypeRuntimeArray %Foo + %A = OpTypeStruct %_runtimearr_Foo +%_ptr_Uniform_A = OpTypePointer Uniform %A + %C1 = OpVariable %_ptr_Uniform_A 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 + %Foo_0 = OpTypeStruct %int %int + %uint_1024 = OpConstant %uint 1024 +%_arr_Foo_0_uint_1024 = OpTypeArray %Foo_0 %uint_1024 + %A_0 = OpTypeStruct %_arr_Foo_0_uint_1024 +%_ptr_Uniform_A_0 = OpTypePointer Uniform %A_0 + %C2 = OpVariable %_ptr_Uniform_A_0 Uniform +%_ptr_Uniform_Foo_0 = OpTypePointer Uniform %Foo_0 +%_ptr_Uniform_Foo = OpTypePointer Uniform %Foo +%_ptr_Uniform_int = OpTypePointer Uniform %int + %int_1 = OpConstant %int 1 +%_runtimearr_Foo_0 = OpTypeRuntimeArray %Foo + %B = OpTypeStruct %_runtimearr_Foo_0 +%_ptr_Uniform_B = OpTypePointer Uniform %B + %C3 = OpVariable %_ptr_Uniform_B Uniform +%_arr_Foo_0_uint_1024_0 = OpTypeArray %Foo_0 %uint_1024 + %B_0 = OpTypeStruct %_arr_Foo_0_uint_1024_0 +%_ptr_Uniform_B_0 = OpTypePointer Uniform %B_0 + %C4 = OpVariable %_ptr_Uniform_B_0 Uniform + %main = OpFunction %void None %3 + %5 = OpLabel + %19 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %20 = OpLoad %uint %19 + %27 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %28 = OpLoad %uint %27 + %30 = OpAccessChain %_ptr_Uniform_Foo_0 %C2 %int_0 %28 + %31 = OpLoad %Foo_0 %30 + %33 = OpAccessChain %_ptr_Uniform_Foo %C1 %int_0 %20 + %34 = OpCompositeExtract %int %31 0 + %36 = OpAccessChain %_ptr_Uniform_int %33 %int_0 + OpStore %36 %34 + %37 = OpCompositeExtract %int %31 1 + %39 = OpAccessChain %_ptr_Uniform_int %33 %int_1 + OpStore %39 %37 + %44 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %45 = OpLoad %uint %44 + %50 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %51 = OpLoad %uint %50 + %52 = OpAccessChain %_ptr_Uniform_Foo_0 %C4 %int_0 %51 + %53 = OpLoad %Foo_0 %52 + %54 = OpAccessChain %_ptr_Uniform_Foo %C3 %int_0 %45 + %55 = OpCompositeExtract %int %53 0 + %56 = OpAccessChain %_ptr_Uniform_int %54 %int_0 + OpStore %56 %55 + %57 = OpCompositeExtract %int %53 1 + %58 = OpAccessChain %_ptr_Uniform_int %54 %int_1 + OpStore %58 %57 + OpReturn + OpFunctionEnd diff --git a/shaders/asm/comp/block-name-alias-global.asm.comp b/shaders/asm/comp/block-name-alias-global.asm.comp new file mode 100644 index 00000000..85f6cc04 --- /dev/null +++ b/shaders/asm/comp/block-name-alias-global.asm.comp @@ -0,0 +1,119 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 59 +; 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 %Foo "A" + OpMemberName %Foo 0 "a" + OpMemberName %Foo 1 "b" + OpName %A "A" + OpMemberName %A 0 "Data" + OpName %C1 "C1" + OpName %gl_GlobalInvocationID "gl_GlobalInvocationID" + OpName %Foo_0 "A" + OpMemberName %Foo_0 0 "a" + OpMemberName %Foo_0 1 "b" + OpName %A_0 "A" + OpMemberName %A_0 0 "Data" + OpName %C2 "C2" + OpName %B "B" + OpMemberName %B 0 "Data" + OpName %C3 "C3" + OpName %B_0 "B" + OpMemberName %B_0 0 "Data" + OpName %C4 "C4" + OpMemberDecorate %Foo 0 Offset 0 + OpMemberDecorate %Foo 1 Offset 4 + OpDecorate %_runtimearr_Foo ArrayStride 8 + OpMemberDecorate %A 0 Offset 0 + OpDecorate %A BufferBlock + OpDecorate %C1 DescriptorSet 0 + OpDecorate %C1 Binding 1 + OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId + OpMemberDecorate %Foo_0 0 Offset 0 + OpMemberDecorate %Foo_0 1 Offset 4 + OpDecorate %_arr_Foo_0_uint_1024 ArrayStride 16 + OpMemberDecorate %A_0 0 Offset 0 + OpDecorate %A_0 Block + OpDecorate %C2 DescriptorSet 0 + OpDecorate %C2 Binding 2 + OpDecorate %_runtimearr_Foo_0 ArrayStride 8 + OpMemberDecorate %B 0 Offset 0 + OpDecorate %B BufferBlock + OpDecorate %C3 DescriptorSet 0 + OpDecorate %C3 Binding 0 + OpDecorate %_arr_Foo_0_uint_1024_0 ArrayStride 16 + OpMemberDecorate %B_0 0 Offset 0 + OpDecorate %B_0 Block + OpDecorate %C4 DescriptorSet 0 + OpDecorate %C4 Binding 3 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %int = OpTypeInt 32 1 + %Foo = OpTypeStruct %int %int +%_runtimearr_Foo = OpTypeRuntimeArray %Foo + %A = OpTypeStruct %_runtimearr_Foo +%_ptr_Uniform_A = OpTypePointer Uniform %A + %C1 = OpVariable %_ptr_Uniform_A 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 + %Foo_0 = OpTypeStruct %int %int + %uint_1024 = OpConstant %uint 1024 +%_arr_Foo_0_uint_1024 = OpTypeArray %Foo_0 %uint_1024 + %A_0 = OpTypeStruct %_arr_Foo_0_uint_1024 +%_ptr_Uniform_A_0 = OpTypePointer Uniform %A_0 + %C2 = OpVariable %_ptr_Uniform_A_0 Uniform +%_ptr_Uniform_Foo_0 = OpTypePointer Uniform %Foo_0 +%_ptr_Uniform_Foo = OpTypePointer Uniform %Foo +%_ptr_Uniform_int = OpTypePointer Uniform %int + %int_1 = OpConstant %int 1 +%_runtimearr_Foo_0 = OpTypeRuntimeArray %Foo + %B = OpTypeStruct %_runtimearr_Foo_0 +%_ptr_Uniform_B = OpTypePointer Uniform %B + %C3 = OpVariable %_ptr_Uniform_B Uniform +%_arr_Foo_0_uint_1024_0 = OpTypeArray %Foo_0 %uint_1024 + %B_0 = OpTypeStruct %_arr_Foo_0_uint_1024_0 +%_ptr_Uniform_B_0 = OpTypePointer Uniform %B_0 + %C4 = OpVariable %_ptr_Uniform_B_0 Uniform + %main = OpFunction %void None %3 + %5 = OpLabel + %19 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %20 = OpLoad %uint %19 + %27 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %28 = OpLoad %uint %27 + %30 = OpAccessChain %_ptr_Uniform_Foo_0 %C2 %int_0 %28 + %31 = OpLoad %Foo_0 %30 + %33 = OpAccessChain %_ptr_Uniform_Foo %C1 %int_0 %20 + %34 = OpCompositeExtract %int %31 0 + %36 = OpAccessChain %_ptr_Uniform_int %33 %int_0 + OpStore %36 %34 + %37 = OpCompositeExtract %int %31 1 + %39 = OpAccessChain %_ptr_Uniform_int %33 %int_1 + OpStore %39 %37 + %44 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %45 = OpLoad %uint %44 + %50 = OpAccessChain %_ptr_Input_uint %gl_GlobalInvocationID %uint_0 + %51 = OpLoad %uint %50 + %52 = OpAccessChain %_ptr_Uniform_Foo_0 %C4 %int_0 %51 + %53 = OpLoad %Foo_0 %52 + %54 = OpAccessChain %_ptr_Uniform_Foo %C3 %int_0 %45 + %55 = OpCompositeExtract %int %53 0 + %56 = OpAccessChain %_ptr_Uniform_int %54 %int_0 + OpStore %56 %55 + %57 = OpCompositeExtract %int %53 1 + %58 = OpAccessChain %_ptr_Uniform_int %54 %int_1 + OpStore %58 %57 + OpReturn + OpFunctionEnd diff --git a/spirv_cross.cpp b/spirv_cross.cpp index 64a85bf2..ee8bb978 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -901,14 +901,30 @@ void Compiler::flatten_interface_block(uint32_t id) var.storage = storage; } -void Compiler::update_name_cache(unordered_set &cache, string &name) +void Compiler::update_name_cache(unordered_set &cache_primary, + const unordered_set &cache_secondary, string &name) { if (name.empty()) return; - if (cache.find(name) == end(cache)) + const auto find_name = [&](const string &n) -> bool { + if (cache_primary.find(n) != end(cache_primary)) + return true; + + if (&cache_primary != &cache_secondary) + if (cache_secondary.find(n) != end(cache_secondary)) + return true; + + return false; + }; + + const auto insert_name = [&](const string &n) { + cache_primary.insert(n); + }; + + if (!find_name(name)) { - cache.insert(name); + insert_name(name); return; } @@ -936,8 +952,13 @@ void Compiler::update_name_cache(unordered_set &cache, string &name) { counter++; name = tmpname + (use_linked_underscore ? "_" : "") + convert_to_string(counter); - } while (cache.find(name) != end(cache)); - cache.insert(name); + } while (find_name(name)); + insert_name(name); +} + +void Compiler::update_name_cache(unordered_set &cache, string &name) +{ + update_name_cache(cache, cache, name); } void Compiler::set_name(uint32_t id, const std::string &name) diff --git a/spirv_cross.hpp b/spirv_cross.hpp index 0767956c..4e2c4bc5 100644 --- a/spirv_cross.hpp +++ b/spirv_cross.hpp @@ -649,6 +649,13 @@ protected: void update_name_cache(std::unordered_set &cache, std::string &name); + // A variant which takes two sets of names. The secondary is only used to verify there are no collisions, + // but the set is not updated when we have found a new name. + // Used primarily when adding block interface names. + void update_name_cache(std::unordered_set &cache_primary, + const std::unordered_set &cache_secondary, + std::string &name); + bool function_is_pure(const SPIRFunction &func); bool block_is_pure(const SPIRBlock &block); bool block_is_outside_flow_control_from_block(const SPIRBlock &from, const SPIRBlock &to); diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index c9428e0f..b77b70a3 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -288,6 +288,7 @@ void CompilerGLSL::reset() block_output_names.clear(); block_ubo_names.clear(); block_ssbo_names.clear(); + block_names.clear(); function_overloads.clear(); for (auto &id : ir.ids) @@ -1544,19 +1545,26 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var) // Shaders never use the block by interface name, so we don't // have to track this other than updating name caches. - if (ir.meta[type.self].decoration.alias.empty() || block_namespace.find(buffer_name) != end(block_namespace)) + // If we have a collision for any reason, just fallback immediately. + if (ir.meta[type.self].decoration.alias.empty() || + block_namespace.find(buffer_name) != end(block_namespace) || + resource_names.find(buffer_name) != end(resource_names)) + { buffer_name = get_block_fallback_name(var.self); + } - // Make sure we get something unique. - add_variable(block_namespace, buffer_name); + // Make sure we get something unique for both global name scope and block name scope. + // See GLSL 4.5 spec: section 4.3.9 for details. + add_variable(block_namespace, resource_names, buffer_name); // If for some reason buffer_name is an illegal name, make a final fallback to a workaround name. // This cannot conflict with anything else, so we're safe now. + // We cannot reuse this fallback name in neither global scope (blocked by block_names) nor block name scope. if (buffer_name.empty()) buffer_name = join("_", get(var.basetype).self, "_", var.self); - // Instance names cannot alias block names. - resource_names.insert(buffer_name); + block_names.insert(buffer_name); + block_namespace.insert(buffer_name); // Save for post-reflection later. declared_block_names[var.self] = buffer_name; @@ -9164,7 +9172,8 @@ string CompilerGLSL::type_to_glsl(const SPIRType &type, uint32_t id) } } -void CompilerGLSL::add_variable(unordered_set &variables, string &name) +void CompilerGLSL::add_variable(unordered_set &variables_primary, + const unordered_set &variables_secondary, string &name) { if (name.empty()) return; @@ -9179,23 +9188,17 @@ void CompilerGLSL::add_variable(unordered_set &variables, string &name) // Avoid double underscores. name = sanitize_underscores(name); - update_name_cache(variables, name); -} - -void CompilerGLSL::add_variable(unordered_set &variables, uint32_t id) -{ - auto &name = ir.meta[id].decoration.alias; - add_variable(variables, name); + update_name_cache(variables_primary, variables_secondary, name); } void CompilerGLSL::add_local_variable_name(uint32_t id) { - add_variable(local_variable_names, id); + add_variable(local_variable_names, block_names, ir.meta[id].decoration.alias); } void CompilerGLSL::add_resource_name(uint32_t id) { - add_variable(resource_names, id); + add_variable(resource_names, block_names, ir.meta[id].decoration.alias); } void CompilerGLSL::add_header_line(const std::string &line) diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 45f7b3f4..d64c8cb3 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -346,6 +346,7 @@ protected: std::unordered_set block_output_names; std::unordered_set block_ubo_names; std::unordered_set block_ssbo_names; + std::unordered_set block_names; // A union of all block_*_names. std::unordered_map> function_overloads; bool processing_entry_point = false; @@ -571,8 +572,13 @@ protected: void emit_pls(); void remap_pls_variables(); - void add_variable(std::unordered_set &variables, uint32_t id); - void add_variable(std::unordered_set &variables, std::string &name); + // A variant which takes two sets of name. The secondary is only used to verify there are no collisions, + // but the set is not updated when we have found a new name. + // Used primarily when adding block interface names. + void add_variable(std::unordered_set &variables_primary, + const std::unordered_set &variables_secondary, + std::string &name); + void check_function_call_constraints(const uint32_t *args, uint32_t length); void handle_invalid_expression(uint32_t id); void find_static_extensions(); diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp index 0979767b..a7df3280 100644 --- a/spirv_hlsl.cpp +++ b/spirv_hlsl.cpp @@ -1893,16 +1893,21 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var) // Prefer the block name if possible. auto buffer_name = to_name(type.self, false); - if (ir.meta[type.self].decoration.alias.empty() || resource_names.find(buffer_name) != end(resource_names)) + if (ir.meta[type.self].decoration.alias.empty() || + resource_names.find(buffer_name) != end(resource_names) || + block_names.find(buffer_name) != end(block_names)) + { buffer_name = get_block_fallback_name(var.self); - add_variable(resource_names, buffer_name); + } + + add_variable(block_names, resource_names, buffer_name); // If for some reason buffer_name is an illegal name, make a final fallback to a workaround name. // This cannot conflict with anything else, so we're safe now. if (buffer_name.empty()) buffer_name = join("_", get(var.basetype).self, "_", var.self); - resource_names.insert(buffer_name); + block_names.insert(buffer_name); // Save for post-reflection later. declared_block_names[var.self] = buffer_name;