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

Flush all variables after storing through a variable pointer.

Browse Source 
Since we can't know which variable was modified, we therefore have to
conservatively assume that any variable might have been modified.
This commit is contained in:
Chip Davis 2019-01-08 15:16:17 -06:00
parent fc02b3d656
commit d6aa911156
5 changed files with 141 additions and 2 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

20
reference/opt/shaders-msl/asm/comp/variable-pointers-store-forwarding.asm.comp Normal file
View File

@ -0,0 +1,20 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct foo
{
int a;
};
struct bar
{
int b;
};
kernel void main0(device foo& x [[buffer(0)]], device bar& y [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
y.b = x.a + x.a;
}

3
reference/shaders-msl-no-opt/asm/comp/variable-pointers.asm.comp
View File

@ -53,7 +53,8 @@ kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(3)]],
_77 = *_73;
if (_77 != 0)
{
int _82 = _77 + (*_76);
int _81 = *_76;
int _82 = _77 + _81;
*_73 = _82;
*_76 = _82;
cur = &_76[1u];

31
reference/shaders-msl/asm/comp/variable-pointers-store-forwarding.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 foo
{
int a;
};
struct bar
{
int b;
};
device int* _24(device foo& a, device bar& b, thread uint3& gl_GlobalInvocationID)
{
return (gl_GlobalInvocationID.x != 0u) ? &a.a : &b.b;
}
kernel void main0(device foo& x [[buffer(0)]], device bar& y [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
device int* _34 = _24(x, y, gl_GlobalInvocationID);
device int* _33 = _34;
int _37 = x.a;
*_33 = 0;
y.b = _37 + _37;
}

75
shaders-msl/asm/comp/variable-pointers-store-forwarding.asm.comp Normal file
View File

@ -0,0 +1,75 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 40
; Schema: 0
OpCapability Shader
OpCapability VariablePointers
%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 "foo"
OpMemberName %foo 0 "a"
OpName %bar "bar"
OpMemberName %bar 0 "b"
OpName %x "x"
OpName %y "y"
OpName %a "a"
OpName %b "b"
OpMemberDecorate %foo 0 Offset 0
OpMemberDecorate %bar 0 Offset 0
OpDecorate %foo Block
OpDecorate %bar Block
OpDecorate %x DescriptorSet 0
OpDecorate %x Binding 0
OpDecorate %y DescriptorSet 0
OpDecorate %y Binding 1
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
%void = OpTypeVoid
%11 = OpTypeFunction %void
%int = OpTypeInt 32 1
%uint = OpTypeInt 32 0
%v3uint = OpTypeVector %uint 3
%_ptr_Input_v3uint = OpTypePointer Input %v3uint
%gl_GlobalInvocationID = OpVariable %_ptr_Input_v3uint Input
%foo = OpTypeStruct %int
%_ptr_StorageBuffer_foo = OpTypePointer StorageBuffer %foo
%x = OpVariable %_ptr_StorageBuffer_foo StorageBuffer
%bar = OpTypeStruct %int
%_ptr_StorageBuffer_bar = OpTypePointer StorageBuffer %bar
%y = OpVariable %_ptr_StorageBuffer_bar StorageBuffer
%uint_0 = OpConstant %uint 0
%int_0 = OpConstant %int 0
%bool = OpTypeBool
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%22 = OpTypeFunction %_ptr_StorageBuffer_int %_ptr_StorageBuffer_foo %_ptr_StorageBuffer_bar
%_ptr_Function__ptr_StorageBuffer_int = OpTypePointer Function %_ptr_StorageBuffer_int
%24 = OpFunction %_ptr_StorageBuffer_int None %22
%a = OpFunctionParameter %_ptr_StorageBuffer_foo
%b = OpFunctionParameter %_ptr_StorageBuffer_bar
%25 = OpLabel
%26 = OpLoad %v3uint %gl_GlobalInvocationID
%27 = OpCompositeExtract %uint %26 0
%28 = OpINotEqual %bool %27 %uint_0
%29 = OpAccessChain %_ptr_StorageBuffer_int %a %uint_0
%30 = OpAccessChain %_ptr_StorageBuffer_int %b %uint_0
%31 = OpSelect %_ptr_StorageBuffer_int %28 %29 %30
OpReturnValue %31
OpFunctionEnd
%main = OpFunction %void None %11
%32 = OpLabel
%33 = OpVariable %_ptr_Function__ptr_StorageBuffer_int Function
%34 = OpFunctionCall %_ptr_StorageBuffer_int %24 %x %y
OpStore %33 %34
%35 = OpLoad %_ptr_StorageBuffer_int %33
%36 = OpAccessChain %_ptr_StorageBuffer_int %x %uint_0
%37 = OpLoad %int %36
OpStore %35 %int_0
%38 = OpIAdd %int %37 %37
%39 = OpAccessChain %_ptr_StorageBuffer_int %y %uint_0
OpStore %39 %38
OpReturn
OpFunctionEnd

14
spirv_cross.cpp
View File

@ -294,7 +294,10 @@ void Compiler::register_write(uint32_t chain)
if (var)
{
// If our variable is in a storage class which can alias with other buffers,
// invalidate all variables which depend on aliased variables.
// invalidate all variables which depend on aliased variables. And if this is a
// variable pointer, then invalidate all variables regardless.
if (get_variable_data_type(*var).pointer)
flush_all_active_variables();
if (variable_storage_is_aliased(*var))
flush_all_aliased_variables();
else if (var)
@ -307,6 +310,15 @@ void Compiler::register_write(uint32_t chain)
force_recompile = true;
}
}
else
{
// If we stored through a variable pointer, then we don't know which
// variable we stored to. So *all* expressions after this point need to
// be invalidated.
// FIXME: If we can prove that the variable pointer will point to
// only certain variables, we can invalidate only those.
flush_all_active_variables();
}
}
void Compiler::flush_dependees(SPIRVariable &var)