AuroraMiddleware/SPIRV-Cross-Vulnerable
1
0
Fork 0
mirror of https://github.com/KhronosGroup/SPIRV-Cross.git synced 2024-11-16 00:40:08 +00:00
Code Issues Projects Releases Wiki Activity

Merge pull request #666 from KhronosGroup/fix-665

Browse Source 
MSL: Fix naming issue of aliased global variables.
This commit is contained in:
Hans-Kristian Arntzen 2018-08-27 10:35:38 +02:00 committed by GitHub
commit d964117cce
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
11 changed files with 449 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -0,0 +1,9 @@
void comp_main()
{
}
[numthreads(1, 1, 1)]
void main()
{
comp_main();
}

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

@ -0,0 +1,9 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
}

7
reference/opt/shaders/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,7 @@
#version 450
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main()
{
}

33
reference/shaders-hlsl/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,33 @@
ByteAddressBuffer ssbo : register(t1);
static uint3 gl_GlobalInvocationID;
struct SPIRV_Cross_Input
{
uint3 gl_GlobalInvocationID : SV_DispatchThreadID;
};
void Load(uint size)
{
int byteAddrTemp = int(size >> uint(2));
uint4 data = uint4(ssbo.Load(byteAddrTemp * 4 + 0), ssbo.Load((byteAddrTemp + 1) * 4 + 0), ssbo.Load((byteAddrTemp + 2) * 4 + 0), ssbo.Load((byteAddrTemp + 3) * 4 + 0));
}
void _main(uint3 id)
{
uint param = 4u;
Load(param);
}
void comp_main()
{
uint3 id = gl_GlobalInvocationID;
uint3 param = id;
_main(param);
}
[numthreads(1, 1, 1)]
void main(SPIRV_Cross_Input stage_input)
{
gl_GlobalInvocationID = stage_input.gl_GlobalInvocationID;
comp_main();
}

31
reference/shaders-msl/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,31 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct ssbo
{
uint _data[1];
};
void Load(thread const uint& size, const device ssbo& ssbo_1)
{
int byteAddrTemp = int(size >> uint(2));
uint4 data = uint4(ssbo_1._data[byteAddrTemp], ssbo_1._data[byteAddrTemp + 1], ssbo_1._data[byteAddrTemp + 2], ssbo_1._data[byteAddrTemp + 3]);
}
void _main(thread const uint3& id, const device ssbo& ssbo_1)
{
uint param = 4u;
Load(param, ssbo_1);
}
kernel void main0(const device ssbo& ssbo_1 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint3 id = gl_GlobalInvocationID;
uint3 param = id;
_main(param, ssbo_1);
}

27
reference/shaders/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,27 @@
#version 450
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
layout(binding = 1, std430) readonly buffer ssbo
{
uint _data[];
} ssbo_1;
void Load(uint size)
{
int byteAddrTemp = int(size >> uint(2));
uvec4 data = uvec4(ssbo_1._data[byteAddrTemp], ssbo_1._data[byteAddrTemp + 1], ssbo_1._data[byteAddrTemp + 2], ssbo_1._data[byteAddrTemp + 3]);
}
void _main(uvec3 id)
{
uint param = 4u;
Load(param);
}
void main()
{
uvec3 id = gl_GlobalInvocationID;
uvec3 param = id;
_main(param);
}

102
shaders-hlsl/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,102 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 6
; Bound: 61
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %id_1
OpExecutionMode %main LocalSize 1 1 1
OpSource HLSL 500
OpName %main "main"
OpName %Load_u1_ "Load(u1;"
OpName %size "size"
OpName %_main_vu3_ "@main(vu3;"
OpName %id "id"
OpName %data "data"
OpName %byteAddrTemp "byteAddrTemp"
OpName %ssbo "ssbo"
OpMemberName %ssbo 0 "@data"
OpName %ssbo_0 "ssbo"
OpName %param "param"
OpName %id_0 "id"
OpName %id_1 "id"
OpName %param_0 "param"
OpDecorate %_runtimearr_uint ArrayStride 4
OpMemberDecorate %ssbo 0 NonWritable
OpMemberDecorate %ssbo 0 Offset 0
OpDecorate %ssbo BufferBlock
OpDecorate %ssbo_0 DescriptorSet 0
OpDecorate %ssbo_0 Binding 1
OpDecorate %id_1 BuiltIn GlobalInvocationId
%void = OpTypeVoid
%3 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint
%8 = OpTypeFunction %void %_ptr_Function_uint
%v3uint = OpTypeVector %uint 3
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%14 = OpTypeFunction %void %_ptr_Function_v3uint
%v4uint = OpTypeVector %uint 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%int_2 = OpConstant %int 2
%_runtimearr_uint = OpTypeRuntimeArray %uint
%ssbo = OpTypeStruct %_runtimearr_uint
%_ptr_Uniform_ssbo = OpTypePointer Uniform %ssbo
%ssbo_0 = OpVariable %_ptr_Uniform_ssbo Uniform
%int_0 = OpConstant %int 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%int_1 = OpConstant %int 1
%int_3 = OpConstant %int 3
%uint_4 = OpConstant %uint 4
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%id_1 = OpVariable %_ptr_Input_v3uint Input
%main = OpFunction %void None %3
%5 = OpLabel
%id_0 = OpVariable %_ptr_Function_v3uint Function
%param_0 = OpVariable %_ptr_Function_v3uint Function
%57 = OpLoad %v3uint %id_1
OpStore %id_0 %57
%59 = OpLoad %v3uint %id_0
OpStore %param_0 %59
%60 = OpFunctionCall %void %_main_vu3_ %param_0
OpReturn
OpFunctionEnd
%Load_u1_ = OpFunction %void None %8
%size = OpFunctionParameter %_ptr_Function_uint
%11 = OpLabel
%data = OpVariable %_ptr_Function_v4uint Function
%byteAddrTemp = OpVariable %_ptr_Function_int Function
%24 = OpLoad %uint %size
%26 = OpShiftRightLogical %int %24 %int_2
OpStore %byteAddrTemp %26
%32 = OpLoad %int %byteAddrTemp
%34 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %32
%35 = OpLoad %uint %34
%36 = OpLoad %int %byteAddrTemp
%38 = OpIAdd %int %36 %int_1
%39 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %38
%40 = OpLoad %uint %39
%41 = OpLoad %int %byteAddrTemp
%42 = OpIAdd %int %41 %int_2
%43 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %42
%44 = OpLoad %uint %43
%45 = OpLoad %int %byteAddrTemp
%47 = OpIAdd %int %45 %int_3
%48 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %47
%49 = OpLoad %uint %48
%50 = OpCompositeConstruct %v4uint %35 %40 %44 %49
OpStore %data %50
OpReturn
OpFunctionEnd
%_main_vu3_ = OpFunction %void None %14
%id = OpFunctionParameter %_ptr_Function_v3uint
%17 = OpLabel
%param = OpVariable %_ptr_Function_uint Function
OpStore %param %uint_4
%53 = OpFunctionCall %void %Load_u1_ %param
OpReturn
OpFunctionEnd

102
shaders-msl/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,102 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 6
; Bound: 61
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %id_1
OpExecutionMode %main LocalSize 1 1 1
OpSource HLSL 500
OpName %main "main"
OpName %Load_u1_ "Load(u1;"
OpName %size "size"
OpName %_main_vu3_ "@main(vu3;"
OpName %id "id"
OpName %data "data"
OpName %byteAddrTemp "byteAddrTemp"
OpName %ssbo "ssbo"
OpMemberName %ssbo 0 "@data"
OpName %ssbo_0 "ssbo"
OpName %param "param"
OpName %id_0 "id"
OpName %id_1 "id"
OpName %param_0 "param"
OpDecorate %_runtimearr_uint ArrayStride 4
OpMemberDecorate %ssbo 0 NonWritable
OpMemberDecorate %ssbo 0 Offset 0
OpDecorate %ssbo BufferBlock
OpDecorate %ssbo_0 DescriptorSet 0
OpDecorate %ssbo_0 Binding 1
OpDecorate %id_1 BuiltIn GlobalInvocationId
%void = OpTypeVoid
%3 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint
%8 = OpTypeFunction %void %_ptr_Function_uint
%v3uint = OpTypeVector %uint 3
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%14 = OpTypeFunction %void %_ptr_Function_v3uint
%v4uint = OpTypeVector %uint 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%int_2 = OpConstant %int 2
%_runtimearr_uint = OpTypeRuntimeArray %uint
%ssbo = OpTypeStruct %_runtimearr_uint
%_ptr_Uniform_ssbo = OpTypePointer Uniform %ssbo
%ssbo_0 = OpVariable %_ptr_Uniform_ssbo Uniform
%int_0 = OpConstant %int 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%int_1 = OpConstant %int 1
%int_3 = OpConstant %int 3
%uint_4 = OpConstant %uint 4
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%id_1 = OpVariable %_ptr_Input_v3uint Input
%main = OpFunction %void None %3
%5 = OpLabel
%id_0 = OpVariable %_ptr_Function_v3uint Function
%param_0 = OpVariable %_ptr_Function_v3uint Function
%57 = OpLoad %v3uint %id_1
OpStore %id_0 %57
%59 = OpLoad %v3uint %id_0
OpStore %param_0 %59
%60 = OpFunctionCall %void %_main_vu3_ %param_0
OpReturn
OpFunctionEnd
%Load_u1_ = OpFunction %void None %8
%size = OpFunctionParameter %_ptr_Function_uint
%11 = OpLabel
%data = OpVariable %_ptr_Function_v4uint Function
%byteAddrTemp = OpVariable %_ptr_Function_int Function
%24 = OpLoad %uint %size
%26 = OpShiftRightLogical %int %24 %int_2
OpStore %byteAddrTemp %26
%32 = OpLoad %int %byteAddrTemp
%34 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %32
%35 = OpLoad %uint %34
%36 = OpLoad %int %byteAddrTemp
%38 = OpIAdd %int %36 %int_1
%39 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %38
%40 = OpLoad %uint %39
%41 = OpLoad %int %byteAddrTemp
%42 = OpIAdd %int %41 %int_2
%43 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %42
%44 = OpLoad %uint %43
%45 = OpLoad %int %byteAddrTemp
%47 = OpIAdd %int %45 %int_3
%48 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %47
%49 = OpLoad %uint %48
%50 = OpCompositeConstruct %v4uint %35 %40 %44 %49
OpStore %data %50
OpReturn
OpFunctionEnd
%_main_vu3_ = OpFunction %void None %14
%id = OpFunctionParameter %_ptr_Function_v3uint
%17 = OpLabel
%param = OpVariable %_ptr_Function_uint Function
OpStore %param %uint_4
%53 = OpFunctionCall %void %Load_u1_ %param
OpReturn
OpFunctionEnd

102
shaders/asm/comp/global-parameter-name-alias.asm.comp Normal file
View File

@ -0,0 +1,102 @@
; SPIR-V
; Version: 1.0
; Generator: Khronos Glslang Reference Front End; 6
; Bound: 61
; Schema: 0
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %id_1
OpExecutionMode %main LocalSize 1 1 1
OpSource HLSL 500
OpName %main "main"
OpName %Load_u1_ "Load(u1;"
OpName %size "size"
OpName %_main_vu3_ "@main(vu3;"
OpName %id "id"
OpName %data "data"
OpName %byteAddrTemp "byteAddrTemp"
OpName %ssbo "ssbo"
OpMemberName %ssbo 0 "@data"
OpName %ssbo_0 "ssbo"
OpName %param "param"
OpName %id_0 "id"
OpName %id_1 "id"
OpName %param_0 "param"
OpDecorate %_runtimearr_uint ArrayStride 4
OpMemberDecorate %ssbo 0 NonWritable
OpMemberDecorate %ssbo 0 Offset 0
OpDecorate %ssbo BufferBlock
OpDecorate %ssbo_0 DescriptorSet 0
OpDecorate %ssbo_0 Binding 1
OpDecorate %id_1 BuiltIn GlobalInvocationId
%void = OpTypeVoid
%3 = OpTypeFunction %void
%uint = OpTypeInt 32 0
%_ptr_Function_uint = OpTypePointer Function %uint
%8 = OpTypeFunction %void %_ptr_Function_uint
%v3uint = OpTypeVector %uint 3
%_ptr_Function_v3uint = OpTypePointer Function %v3uint
%14 = OpTypeFunction %void %_ptr_Function_v3uint
%v4uint = OpTypeVector %uint 4
%_ptr_Function_v4uint = OpTypePointer Function %v4uint
%int = OpTypeInt 32 1
%_ptr_Function_int = OpTypePointer Function %int
%int_2 = OpConstant %int 2
%_runtimearr_uint = OpTypeRuntimeArray %uint
%ssbo = OpTypeStruct %_runtimearr_uint
%_ptr_Uniform_ssbo = OpTypePointer Uniform %ssbo
%ssbo_0 = OpVariable %_ptr_Uniform_ssbo Uniform
%int_0 = OpConstant %int 0
%_ptr_Uniform_uint = OpTypePointer Uniform %uint
%int_1 = OpConstant %int 1
%int_3 = OpConstant %int 3
%uint_4 = OpConstant %uint 4
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%id_1 = OpVariable %_ptr_Input_v3uint Input
%main = OpFunction %void None %3
%5 = OpLabel
%id_0 = OpVariable %_ptr_Function_v3uint Function
%param_0 = OpVariable %_ptr_Function_v3uint Function
%57 = OpLoad %v3uint %id_1
OpStore %id_0 %57
%59 = OpLoad %v3uint %id_0
OpStore %param_0 %59
%60 = OpFunctionCall %void %_main_vu3_ %param_0
OpReturn
OpFunctionEnd
%Load_u1_ = OpFunction %void None %8
%size = OpFunctionParameter %_ptr_Function_uint
%11 = OpLabel
%data = OpVariable %_ptr_Function_v4uint Function
%byteAddrTemp = OpVariable %_ptr_Function_int Function
%24 = OpLoad %uint %size
%26 = OpShiftRightLogical %int %24 %int_2
OpStore %byteAddrTemp %26
%32 = OpLoad %int %byteAddrTemp
%34 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %32
%35 = OpLoad %uint %34
%36 = OpLoad %int %byteAddrTemp
%38 = OpIAdd %int %36 %int_1
%39 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %38
%40 = OpLoad %uint %39
%41 = OpLoad %int %byteAddrTemp
%42 = OpIAdd %int %41 %int_2
%43 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %42
%44 = OpLoad %uint %43
%45 = OpLoad %int %byteAddrTemp
%47 = OpIAdd %int %45 %int_3
%48 = OpAccessChain %_ptr_Uniform_uint %ssbo_0 %int_0 %47
%49 = OpLoad %uint %48
%50 = OpCompositeConstruct %v4uint %35 %40 %44 %49
OpStore %data %50
OpReturn
OpFunctionEnd
%_main_vu3_ = OpFunction %void None %14
%id = OpFunctionParameter %_ptr_Function_v3uint
%17 = OpLabel
%param = OpVariable %_ptr_Function_uint Function
OpStore %param %uint_4
%53 = OpFunctionCall %void %Load_u1_ %param
OpReturn
OpFunctionEnd

10
spirv_glsl.cpp
View File

@ -2271,7 +2271,12 @@ void CompilerGLSL::emit_resources()
// Subclasses may override to modify the return value.
string CompilerGLSL::to_func_call_arg(uint32_t id)
{
return to_expression(id);
// Make sure that we use the name of the original variable, and not the parameter alias.
uint32_t name_id = id;
auto *var = maybe_get<SPIRVariable>(id);
if (var && var->basevariable)
name_id = var->basevariable;
return to_expression(name_id);
}
void CompilerGLSL::handle_invalid_expression(uint32_t id)
@ -8104,13 +8109,14 @@ void CompilerGLSL::append_global_func_args(const SPIRFunction &func, uint32_t in
{
auto &arg = args[arg_idx];
assert(arg.alias_global_variable);
arglist.push_back(to_func_call_arg(arg.id));
// If the underlying variable needs to be declared
// (ie. a local variable with deferred declaration), do so now.
uint32_t var_id = get<SPIRVariable>(arg.id).basevariable;
if (var_id)
flush_variable_declaration(var_id);
arglist.push_back(to_func_call_arg(arg.id));
}
}

24
spirv_msl.cpp
View File

@ -2408,16 +2408,23 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &)
for (auto &arg : func.arguments)
{
add_local_variable_name(arg.id);
uint32_t name_id = arg.id;
string address_space;
auto *var = maybe_get<SPIRVariable>(arg.id);
if (var)
{
// If we need to modify the name of the variable, make sure we modify the original variable.
// Our alias is just a shadow variable.
if (arg.alias_global_variable && var->basevariable)
name_id = var->basevariable;
var->parameter = &arg; // Hold a pointer to the parameter so we can invalidate the readonly field if needed.
address_space = get_argument_address_space(*var);
}
add_local_variable_name(name_id);
if (!address_space.empty())
decl += address_space + " ";
decl += argument_decl(arg);
@ -3408,9 +3415,16 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base
string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
{
auto &var = get<SPIRVariable>(arg.id);
auto &type = expression_type(arg.id);
// If we need to modify the name of the variable, make sure we use the original variable.
// Our alias is just a shadow variable.
uint32_t name_id = var.self;
if (arg.alias_global_variable && var.basevariable)
name_id = var.basevariable;
bool constref = !arg.alias_global_variable && type.pointer && arg.write_count == 0;
bool type_is_image = type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage ||
@ -3437,13 +3451,13 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
{
// If the argument is a pure value and not an opaque type, we will pass by value.
decl += " ";
decl += to_expression(var.self);
decl += to_expression(name_id);
}
else if (is_array(type) && !type_is_image)
{
// Arrays of images and samplers are special cased.
decl += " (&";
decl += to_expression(var.self);
decl += to_expression(name_id);
decl += ")";
decl += type_to_array_glsl(type);
}
@ -3451,12 +3465,12 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
{
decl += "&";
decl += " ";
decl += to_expression(var.self);
decl += to_expression(name_id);
}
else
{
decl += " ";
decl += to_expression(var.self);
decl += to_expression(name_id);
}
return decl;