AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Use correct block-name / other-name aliasing rules.

Browse Source 
A block name cannot alias with any name in its own scope,
and it cannot alias with any other "global" name.

To solve this, we need to complicate the name cache updates a little bit
where we have a "primary" namespace and "secondary" namespace.
This commit is contained in:
Hans-Kristian Arntzen 2019-01-04 12:15:43 +01:00
parent ed16b3e699
commit 9728f9c1b7
14 changed files with 690 additions and 25 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

45
reference/opt/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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();
}

45
reference/opt/shaders-msl/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -0,0 +1,45 @@
#include <metal_stdlib>
#include <simd/simd.h>
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;
}

43
reference/opt/shaders/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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;
}

45
reference/shaders-hlsl/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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();
}

45
reference/shaders-msl/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -0,0 +1,45 @@
#include <metal_stdlib>
#include <simd/simd.h>
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;
}

43
reference/shaders/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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;
}

119
shaders-hlsl/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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

119
shaders-msl/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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

119
shaders/asm/comp/block-name-alias-global.asm.comp Normal file
View File

@ -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

31
spirv_cross.cpp
View File

@ -901,14 +901,30 @@ void Compiler::flatten_interface_block(uint32_t id)
var.storage = storage;
}
void Compiler::update_name_cache(unordered_set<string> &cache, string &name)
void Compiler::update_name_cache(unordered_set<string> &cache_primary,
const unordered_set<string> &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<string> &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<string> &cache, string &name)
{
update_name_cache(cache, cache, name);
}
void Compiler::set_name(uint32_t id, const std::string &name)

7
spirv_cross.hpp
View File

@ -649,6 +649,13 @@ protected:
void update_name_cache(std::unordered_set<std::string> &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<std::string> &cache_primary,
const std::unordered_set<std::string> &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);

33
spirv_glsl.cpp
View File

@ -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<SPIRType>(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<string> &variables, string &name)
void CompilerGLSL::add_variable(unordered_set<string> &variables_primary,
const unordered_set<string> &variables_secondary, string &name)
{
if (name.empty())
return;
@ -9179,23 +9188,17 @@ void CompilerGLSL::add_variable(unordered_set<string> &variables, string &name)
// Avoid double underscores.
name = sanitize_underscores(name);
update_name_cache(variables, name);
}
void CompilerGLSL::add_variable(unordered_set<string> &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)

10
spirv_glsl.hpp
View File

@ -346,6 +346,7 @@ protected:
std::unordered_set<std::string> block_output_names;
std::unordered_set<std::string> block_ubo_names;
std::unordered_set<std::string> block_ssbo_names;
std::unordered_set<std::string> block_names; // A union of all block_*_names.
std::unordered_map<std::string, std::unordered_set<uint64_t>> 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<std::string> &variables, uint32_t id);
void add_variable(std::unordered_set<std::string> &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<std::string> &variables_primary,
const std::unordered_set<std::string> &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();

11
spirv_hlsl.cpp
View File

@ -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<SPIRType>(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;