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

MSL: Support SPV_KHR_variable_pointers.

Browse Source 
This allows shaders to declare and use pointer-type variables. Pointers
may be loaded and stored, be the result of an `OpSelect`, be passed to
and returned from functions, and even be passed as inputs to the `OpPhi`
instruction. All types of pointers may be used as variable pointers.
Variable pointers to storage buffers and workgroup memory may even be
loaded from and stored to, as though they were ordinary variables. In
addition, this enables using an interior pointer to an array as though
it were an array pointer itself using the `OpPtrAccessChain`
instruction.

This is a rather large and involved change, mostly because this is
somewhat complicated with a lot of moving parts. It's a wonder
SPIRV-Cross's output is largely unchanged. Indeed, many of these changes
are to accomplish exactly that! Perhaps the largest source of changes
was the violation of the assumption that, when emitting types, the
pointer type didn't matter.

One of the test cases added by the change doesn't optimize very well;
the output of `spirv-opt` here is invalid SPIR-V. I need to file a bug
with SPIRV-Tools about this.

I wanted to test that variable pointers to images worked too, but I
couldn't figure out how to propagate the access qualifier properly--in
MSL, it's part of the type, so getting this right is important. I've
punted on that for now.
This commit is contained in:
Chip Davis 2018-12-03 02:06:33 -06:00
parent b30d21bafd
commit 3bfb2f94d4
13 changed files with 855 additions and 83 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

31
reference/opt/shaders-msl/asm/comp/variable-pointers-2.asm.comp Normal file
View File

@ -0,0 +1,31 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct foo
{
int a[128];
uint b;
float2 c;
};
struct bar
{
int d;
};
kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{
bool _70 = cb.d != 0;
for (device int* _52 = &(_70 ? &buf : nullptr)->a[0u], * _55 = &buf.a[0u]; (*_52) != (*_55); )
{
int _66 = ((*_52) + (*_55)) + int((*(_70 ? &gl_GlobalInvocationID : &gl_LocalInvocationID)).x);
*_52 = _66;
*_55 = _66;
_52 = &_52[1u];
_55 = &_55[1u];
continue;
}
}

69
reference/shaders-msl-no-opt/asm/comp/variable-pointers.asm.comp Normal file
View File

@ -0,0 +1,69 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct foo
{
int a[128];
uint b;
float2 c;
};
struct bar
{
int d;
};
struct baz
{
int e[128];
};
device int* select_buffer(device foo& buf, device baz& buf2, constant bar& cb)
{
return (cb.d != 0) ? &buf.a[0u] : &buf2.e[0u];
}
device int* select_buffer_null(device foo& buf, constant bar& cb)
{
return (cb.d != 0) ? &buf.a[0u] : nullptr;
}
threadgroup int* select_tgsm(constant bar& cb, threadgroup int (&tgsm)[128])
{
return (cb.d != 0) ? &tgsm[0u] : nullptr;
}
kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(3)]], device baz& buf2 [[buffer(4)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
threadgroup int tgsm[128];
device int* sbuf = select_buffer(buf, buf2, cb);
device int* sbuf2 = select_buffer_null(buf, cb);
threadgroup int* stgsm = select_tgsm(cb, tgsm);
threadgroup int* cur = stgsm;
device int* _73;
_73 = &buf.a[0u];
int _77;
for (;;)
{
threadgroup int* _76 = cur;
_77 = *_73;
if (_77 != 0)
{
int _82 = _77 + (*_76);
*_73 = _82;
*_76 = _82;
cur = &_76[1u];
_73 = &_73[1u];
continue;
}
else
{
break;
}
}
}

42
reference/shaders-msl/asm/comp/variable-pointers-2.asm.comp Normal file
View File

@ -0,0 +1,42 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct foo
{
int a[128];
uint b;
float2 c;
};
struct bar
{
int d;
};
device foo* select_buffer(device foo& a, constant bar& cb)
{
return (cb.d != 0) ? &a : nullptr;
}
thread uint3* select_input(thread uint3& gl_GlobalInvocationID, thread uint3& gl_LocalInvocationID, constant bar& cb)
{
return (cb.d != 0) ? &gl_GlobalInvocationID : &gl_LocalInvocationID;
}
kernel void main0(device foo& buf [[buffer(0)]], constant bar& cb [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint3 gl_LocalInvocationID [[thread_position_in_threadgroup]])
{
device foo* _46 = select_buffer(buf, cb);
device foo* _45 = _46;
for (device int* _52 = &_45->a[0u], * _55 = &buf.a[0u]; (*_52) != (*_55); _52 = &_52[1u], _55 = &_55[1u])
{
int _66 = ((*_52) + (*_55)) + int((*select_input(gl_GlobalInvocationID, gl_LocalInvocationID, cb)).x);
*_52 = _66;
*_55 = _66;
continue;
}
}

152
shaders-msl-no-opt/asm/comp/variable-pointers.asm.comp Normal file
View File

@ -0,0 +1,152 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 89
; 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"
OpMemberName %foo 1 "b"
OpMemberName %foo 2 "c"
OpName %bar "bar"
OpMemberName %bar 0 "d"
OpName %baz "baz"
OpMemberName %baz 0 "e"
OpName %buf "buf"
OpName %buf2 "buf2"
OpName %cb "cb"
OpName %tgsm "tgsm"
OpName %sbuf "sbuf"
OpName %sbuf2 "sbuf2"
OpName %stgsm "stgsm"
OpName %select_buffer "select_buffer"
OpName %select_buffer_null "select_buffer_null"
OpName %select_tgsm "select_tgsm"
OpName %cur "cur"
OpMemberDecorate %foo 0 Offset 0
OpMemberDecorate %foo 1 Offset 512
OpMemberDecorate %foo 2 Offset 520
OpMemberDecorate %bar 0 Offset 0
OpMemberDecorate %baz 0 Offset 0
OpDecorate %foo Block
OpDecorate %bar Block
OpDecorate %baz Block
OpDecorate %buf DescriptorSet 0
OpDecorate %buf Binding 0
OpDecorate %cb DescriptorSet 0
OpDecorate %cb Binding 3
OpDecorate %buf2 DescriptorSet 0
OpDecorate %buf2 Binding 4
OpDecorate %_ptr_Workgroup_int ArrayStride 4
OpDecorate %_ptr_StorageBuffer_int ArrayStride 4
OpDecorate %_arr_int_uint_128 ArrayStride 4
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
%void = OpTypeVoid
%22 = 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
%uint_128 = OpConstant %uint 128
%_arr_int_uint_128 = OpTypeArray %int %uint_128
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%foo = OpTypeStruct %_arr_int_uint_128 %uint %v2float
%_ptr_StorageBuffer_foo = OpTypePointer StorageBuffer %foo
%buf = OpVariable %_ptr_StorageBuffer_foo StorageBuffer
%bar = OpTypeStruct %int
%_ptr_Uniform_bar = OpTypePointer Uniform %bar
%cb = OpVariable %_ptr_Uniform_bar Uniform
%baz = OpTypeStruct %_arr_int_uint_128
%_ptr_StorageBuffer_baz = OpTypePointer StorageBuffer %baz
%buf2 = OpVariable %_ptr_StorageBuffer_baz StorageBuffer
%_ptr_Workgroup__arr_int_uint_128 = OpTypePointer Workgroup %_arr_int_uint_128
%tgsm = OpVariable %_ptr_Workgroup__arr_int_uint_128 Workgroup
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%_ptr_Private__ptr_StorageBuffer_int = OpTypePointer Private %_ptr_StorageBuffer_int
%sbuf = OpVariable %_ptr_Private__ptr_StorageBuffer_int Private
%sbuf2 = OpVariable %_ptr_Private__ptr_StorageBuffer_int Private
%_ptr_Workgroup_int = OpTypePointer Workgroup %int
%_ptr_Private__ptr_Workgroup_int = OpTypePointer Private %_ptr_Workgroup_int
%stgsm = OpVariable %_ptr_Private__ptr_Workgroup_int Private
%uint_0 = OpConstant %uint 0
%bool = OpTypeBool
%_ptr_Uniform_int = OpTypePointer Uniform %int
%44 = OpTypeFunction %_ptr_StorageBuffer_int
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%47 = OpConstantNull %_ptr_StorageBuffer_int
%48 = OpTypeFunction %_ptr_Workgroup_int
%49 = OpConstantNull %_ptr_Workgroup_int
%_ptr_Function__ptr_Workgroup_int = OpTypePointer Function %_ptr_Workgroup_int
%select_buffer = OpFunction %_ptr_StorageBuffer_int None %44
%51 = OpLabel
%52 = OpAccessChain %_ptr_Uniform_int %cb %uint_0
%53 = OpLoad %int %52
%54 = OpINotEqual %bool %53 %int_0
%55 = OpAccessChain %_ptr_StorageBuffer_int %buf %uint_0 %uint_0
%56 = OpAccessChain %_ptr_StorageBuffer_int %buf2 %uint_0 %uint_0
%57 = OpSelect %_ptr_StorageBuffer_int %54 %55 %56
OpReturnValue %57
OpFunctionEnd
%select_buffer_null = OpFunction %_ptr_StorageBuffer_int None %44
%58 = OpLabel
%59 = OpAccessChain %_ptr_Uniform_int %cb %uint_0
%60 = OpLoad %int %59
%61 = OpINotEqual %bool %60 %int_0
%62 = OpAccessChain %_ptr_StorageBuffer_int %buf %uint_0 %uint_0
%63 = OpSelect %_ptr_StorageBuffer_int %61 %62 %47
OpReturnValue %63
OpFunctionEnd
%select_tgsm = OpFunction %_ptr_Workgroup_int None %48
%64 = OpLabel
%65 = OpAccessChain %_ptr_Uniform_int %cb %uint_0
%66 = OpLoad %int %65
%67 = OpINotEqual %bool %66 %int_0
%68 = OpAccessChain %_ptr_Workgroup_int %tgsm %uint_0
%69 = OpSelect %_ptr_Workgroup_int %67 %68 %49
OpReturnValue %69
OpFunctionEnd
%main = OpFunction %void None %22
%70 = OpLabel
%cur = OpVariable %_ptr_Function__ptr_Workgroup_int Function
%71 = OpFunctionCall %_ptr_StorageBuffer_int %select_buffer
OpStore %sbuf %71
%72 = OpFunctionCall %_ptr_StorageBuffer_int %select_buffer_null
OpStore %sbuf2 %72
%73 = OpFunctionCall %_ptr_Workgroup_int %select_tgsm
OpStore %stgsm %73
%74 = OpAccessChain %_ptr_StorageBuffer_int %buf %uint_0 %uint_0
%75 = OpLoad %_ptr_Workgroup_int %stgsm
%76 = OpCopyObject %_ptr_Workgroup_int %75
OpStore %cur %76
OpBranch %77
%77 = OpLabel
%78 = OpPhi %_ptr_StorageBuffer_int %74 %70 %79 %80
%81 = OpLoad %_ptr_Workgroup_int %cur
%82 = OpLoad %int %78
%83 = OpINotEqual %bool %82 %int_0
OpLoopMerge %85 %80 None
OpBranchConditional %83 %84 %85
%84 = OpLabel
%86 = OpLoad %int %81
%87 = OpIAdd %int %82 %86
OpStore %78 %87
OpStore %81 %87
OpBranch %80
%80 = OpLabel
%79 = OpPtrAccessChain %_ptr_StorageBuffer_int %78 %uint_1
%88 = OpPtrAccessChain %_ptr_Workgroup_int %81 %uint_1
OpStore %cur %88
OpBranch %77
%85 = OpLabel
OpReturn
OpFunctionEnd

117
shaders-msl/asm/comp/variable-pointers-2.asm.comp Normal file
View File

@ -0,0 +1,117 @@
; SPIR-V
; Version: 1.3
; Generator: Khronos SPIR-V Tools Assembler; 0
; Bound: 65
; Schema: 0
OpCapability Shader
OpCapability VariablePointers
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main" %gl_GlobalInvocationID %gl_LocalInvocationID
OpExecutionMode %main LocalSize 1 1 1
OpSource GLSL 450
OpName %main "main"
OpName %foo "foo"
OpMemberName %foo 0 "a"
OpMemberName %foo 1 "b"
OpMemberName %foo 2 "c"
OpName %bar "bar"
OpMemberName %bar 0 "d"
OpName %buf "buf"
OpName %cb "cb"
OpName %select_buffer "select_buffer"
OpName %select_input "select_input"
OpName %a "a"
OpMemberDecorate %foo 0 Offset 0
OpMemberDecorate %foo 1 Offset 512
OpMemberDecorate %foo 2 Offset 520
OpMemberDecorate %bar 0 Offset 0
OpDecorate %foo Block
OpDecorate %bar Block
OpDecorate %buf DescriptorSet 0
OpDecorate %buf Binding 0
OpDecorate %cb DescriptorSet 0
OpDecorate %cb Binding 1
OpDecorate %_ptr_StorageBuffer_int ArrayStride 4
OpDecorate %_arr_int_uint_128 ArrayStride 4
OpDecorate %gl_GlobalInvocationID BuiltIn GlobalInvocationId
OpDecorate %gl_LocalInvocationID BuiltIn LocalInvocationId
%void = OpTypeVoid
%15 = 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
%gl_LocalInvocationID = OpVariable %_ptr_Input_v3uint Input
%uint_128 = OpConstant %uint 128
%_arr_int_uint_128 = OpTypeArray %int %uint_128
%float = OpTypeFloat 32
%v2float = OpTypeVector %float 2
%foo = OpTypeStruct %_arr_int_uint_128 %uint %v2float
%_ptr_StorageBuffer_foo = OpTypePointer StorageBuffer %foo
%buf = OpVariable %_ptr_StorageBuffer_foo StorageBuffer
%bar = OpTypeStruct %int
%_ptr_Uniform_bar = OpTypePointer Uniform %bar
%cb = OpVariable %_ptr_Uniform_bar Uniform
%uint_0 = OpConstant %uint 0
%bool = OpTypeBool
%_ptr_Uniform_int = OpTypePointer Uniform %int
%28 = OpTypeFunction %_ptr_StorageBuffer_foo %_ptr_StorageBuffer_foo
%int_0 = OpConstant %int 0
%uint_1 = OpConstant %uint 1
%31 = OpConstantNull %_ptr_StorageBuffer_foo
%32 = OpTypeFunction %_ptr_Input_v3uint
%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
%_ptr_Function__ptr_StorageBuffer_foo = OpTypePointer Function %_ptr_StorageBuffer_foo
%select_buffer = OpFunction %_ptr_StorageBuffer_foo None %28
%a = OpFunctionParameter %_ptr_StorageBuffer_foo
%33 = OpLabel
%34 = OpAccessChain %_ptr_Uniform_int %cb %uint_0
%35 = OpLoad %int %34
%36 = OpINotEqual %bool %35 %int_0
%37 = OpSelect %_ptr_StorageBuffer_foo %36 %a %31
OpReturnValue %37
OpFunctionEnd
%select_input = OpFunction %_ptr_Input_v3uint None %32
%38 = OpLabel
%39 = OpAccessChain %_ptr_Uniform_int %cb %uint_0
%40 = OpLoad %int %39
%41 = OpINotEqual %bool %40 %int_0
%42 = OpSelect %_ptr_Input_v3uint %41 %gl_GlobalInvocationID %gl_LocalInvocationID
OpReturnValue %42
OpFunctionEnd
%main = OpFunction %void None %15
%43 = OpLabel
%65 = OpVariable %_ptr_Function__ptr_StorageBuffer_foo Function
%44 = OpFunctionCall %_ptr_StorageBuffer_foo %select_buffer %buf
OpStore %65 %44
%45 = OpFunctionCall %_ptr_Input_v3uint %select_input
%66 = OpLoad %_ptr_StorageBuffer_foo %65
%46 = OpAccessChain %_ptr_StorageBuffer_int %66 %uint_0 %uint_0
%47 = OpAccessChain %_ptr_StorageBuffer_int %buf %uint_0 %uint_0
OpBranch %48
%48 = OpLabel
%49 = OpPhi %_ptr_StorageBuffer_int %46 %43 %50 %51
%52 = OpPhi %_ptr_StorageBuffer_int %47 %43 %53 %51
%54 = OpLoad %int %49
%55 = OpLoad %int %52
%56 = OpINotEqual %bool %54 %55
OpLoopMerge %58 %51 None
OpBranchConditional %56 %57 %58
%57 = OpLabel
%59 = OpIAdd %int %54 %55
%60 = OpLoad %v3uint %45
%61 = OpCompositeExtract %uint %60 0
%62 = OpBitcast %int %61
%63 = OpIAdd %int %59 %62
OpStore %49 %63
OpStore %52 %63
OpBranch %51
%51 = OpLabel
%50 = OpPtrAccessChain %_ptr_StorageBuffer_int %49 %uint_1
%53 = OpPtrAccessChain %_ptr_StorageBuffer_int %52 %uint_1
OpBranch %48
%58 = OpLabel
OpReturn
OpFunctionEnd

3
spirv_common.hpp
View File

@ -564,6 +564,9 @@ struct SPIRExpression : IVariant
// This is needed for targets which don't support row_major layouts.
bool need_transpose = false;
// Whether or not this is an access chain expression.
bool access_chain = false;
// A list of expressions which this expression depends on.
std::vector<uint32_t> expression_dependencies;

111
spirv_cross.cpp
View File

@ -546,6 +546,40 @@ bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t
break;
}
case OpSelect:
{
// Invalid SPIR-V.
if (length < 5)
return false;
uint32_t count = length - 3;
args += 3;
for (uint32_t i = 0; i < count; i++)
{
auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
if (var && storage_class_is_interface(var->storage))
variables.insert(args[i]);
}
break;
}
case OpPhi:
{
// Invalid SPIR-V.
if (length < 2)
return false;
uint32_t count = length - 2;
args += 2;
for (uint32_t i = 0; i < count; i += 2)
{
auto *var = compiler.maybe_get<SPIRVariable>(args[i]);
if (var && storage_class_is_interface(var->storage))
variables.insert(args[i]);
}
break;
}
case OpAtomicStore:
case OpStore:
// Invalid SPIR-V.
@ -602,6 +636,7 @@ bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t
case OpAccessChain:
case OpInBoundsAccessChain:
case OpPtrAccessChain:
case OpLoad:
case OpCopyObject:
case OpImageTexelPointer:
@ -898,6 +933,7 @@ void Compiler::flatten_interface_block(uint32_t id)
type.array.push_back(array_size);
type.pointer = true;
type.storage = storage;
type.parent_type = t;
var.storage = storage;
}
@ -977,11 +1013,10 @@ const SPIRType &Compiler::get_type_from_variable(uint32_t id) const
uint32_t Compiler::get_non_pointer_type_id(uint32_t type_id) const
{
auto *p_type = &get<SPIRType>(type_id);
while (p_type->pointer)
if (p_type->pointer)
{
assert(p_type->parent_type);
type_id = p_type->parent_type;
p_type = &get<SPIRType>(type_id);
}
return type_id;
}
@ -989,7 +1024,7 @@ uint32_t Compiler::get_non_pointer_type_id(uint32_t type_id) const
const SPIRType &Compiler::get_non_pointer_type(const SPIRType &type) const
{
auto *p_type = &type;
while (p_type->pointer)
if (p_type->pointer)
{
assert(p_type->parent_type);
p_type = &get<SPIRType>(p_type->parent_type);
@ -1002,6 +1037,23 @@ const SPIRType &Compiler::get_non_pointer_type(uint32_t type_id) const
return get_non_pointer_type(get<SPIRType>(type_id));
}
uint32_t Compiler::get_variable_data_type_id(const SPIRVariable &var) const
{
if (var.phi_variable)
return var.basetype;
return get_non_pointer_type_id(var.basetype);
}
SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var)
{
return get<SPIRType>(get_variable_data_type_id(var));
}
const SPIRType &Compiler::get_variable_data_type(const SPIRVariable &var) const
{
return get<SPIRType>(get_variable_data_type_id(var));
}
bool Compiler::is_sampled_image_type(const SPIRType &type)
{
return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 &&
@ -1487,11 +1539,13 @@ size_t Compiler::get_declared_struct_member_size(const SPIRType &struct_type, ui
bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint32_t length)
{
if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain)
if (opcode != OpAccessChain && opcode != OpInBoundsAccessChain && opcode != OpPtrAccessChain)
return true;
bool ptr_chain = (opcode == OpPtrAccessChain);
// Invalid SPIR-V.
if (length < 4)
if (length < (ptr_chain ? 5 : 4))
return false;
if (args[2] != id)
@ -1499,7 +1553,7 @@ bool Compiler::BufferAccessHandler::handle(Op opcode, const uint32_t *args, uint
// Don't bother traversing the entire access chain tree yet.
// If we access a struct member, assume we access the entire member.
uint32_t index = compiler.get<SPIRConstant>(args[3]).scalar();
uint32_t index = compiler.get<SPIRConstant>(args[ptr_chain ? 4 : 3]).scalar();
// Seen this index already.
if (seen.find(index) != end(seen))
@ -2043,6 +2097,7 @@ void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIR
ptr_type = type;
ptr_type.pointer = true;
ptr_type.storage = StorageClassUniformConstant;
ptr_type.parent_type = type_id;
// Build new variable.
compiler.set<SPIRVariable>(combined_id, ptr_type_id, StorageClassFunction, 0);
@ -2116,6 +2171,7 @@ bool Compiler::DummySamplerForCombinedImageHandler::handle(Op opcode, const uint
case OpInBoundsAccessChain:
case OpAccessChain:
case OpPtrAccessChain:
{
if (length < 3)
return false;
@ -2175,6 +2231,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar
case OpInBoundsAccessChain:
case OpAccessChain:
case OpPtrAccessChain:
{
if (length < 3)
return false;
@ -2303,6 +2360,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar
type = base;
type.pointer = true;
type.storage = StorageClassUniformConstant;
type.parent_type = type_id;
// Build new variable.
compiler.set<SPIRVariable>(combined_id, type_id, StorageClassUniformConstant, 0);
@ -2350,6 +2408,7 @@ uint32_t Compiler::build_dummy_sampler_for_combined_images()
ptr_sampler.self = type_id;
ptr_sampler.storage = StorageClassUniformConstant;
ptr_sampler.pointer = true;
ptr_sampler.parent_type = type_id;
set<SPIRVariable>(var_id, ptr_type_id, StorageClassUniformConstant, 0);
set_name(var_id, "SPIRV_Cross_DummySampler");
@ -2599,6 +2658,7 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3
case OpAccessChain:
case OpInBoundsAccessChain:
case OpPtrAccessChain:
{
if (length < 3)
return false;
@ -2783,6 +2843,7 @@ bool Compiler::StaticExpressionAccessHandler::handle(spv::Op op, const uint32_t
case OpAccessChain:
case OpInBoundsAccessChain:
case OpPtrAccessChain:
if (length < 3)
return false;
if (args[2] == variable_id) // If we try to access chain our candidate variable before we store to it, bail.
@ -3219,6 +3280,26 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args
add_if_builtin(args[2]);
break;
case OpSelect:
if (length < 5)
return false;
add_if_builtin(args[3]);
add_if_builtin(args[4]);
break;
case OpPhi:
{
if (length < 2)
return false;
uint32_t count = length - 2;
args += 2;
for (uint32_t i = 0; i < count; i += 2)
add_if_builtin(args[i]);
break;
}
case OpFunctionCall:
{
if (length < 3)
@ -3233,6 +3314,7 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args
case OpAccessChain:
case OpInBoundsAccessChain:
case OpPtrAccessChain:
{
if (length < 4)
return false;
@ -3247,7 +3329,7 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args
add_if_builtin(args[2]);
// Start traversing type hierarchy at the proper non-pointer types.
auto *type = &compiler.get_non_pointer_type(var->basetype);
auto *type = &compiler.get_variable_data_type(*var);
auto &flags =
type->storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins;
@ -3256,6 +3338,13 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args
args += 3;
for (uint32_t i = 0; i < count; i++)
{
// Pointers
if (opcode == OpPtrAccessChain && i == 0)
{
type = &compiler.get<SPIRType>(type->parent_type);
continue;
}
// Arrays
if (!type->array.empty())
{
@ -3463,6 +3552,7 @@ bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_
{
case OpAccessChain:
case OpInBoundsAccessChain:
case OpPtrAccessChain:
case OpLoad:
{
if (length < 3)
@ -3533,7 +3623,12 @@ void Compiler::make_constant_null(uint32_t id, uint32_t type)
{
auto &constant_type = get<SPIRType>(type);
if (!constant_type.array.empty())
if (constant_type.pointer)
{
auto &constant = set<SPIRConstant>(id, type);
constant.make_null(constant_type);
}
else if (!constant_type.array.empty())
{
assert(constant_type.parent_type);
uint32_t parent_id = ir.increase_bound_by(1);

9
spirv_cross.hpp
View File

@ -186,6 +186,15 @@ public:
// Gets the SPIR-V type underlying the given type_id, which might be a pointer.
const SPIRType &get_non_pointer_type(uint32_t type_id) const;
// Gets the ID of the SPIR-V type underlying a variable.
uint32_t get_variable_data_type_id(const SPIRVariable &var) const;
// Gets the SPIR-V type underlying a variable.
SPIRType &get_variable_data_type(const SPIRVariable &var);
// Gets the SPIR-V type underlying a variable.
const SPIRType &get_variable_data_type(const SPIRVariable &var) const;
// Returns if the given type refers to a sampled image.
bool is_sampled_image_type(const SPIRType &type);

255
spirv_glsl.cpp
View File

@ -1756,7 +1756,8 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var)
else
{
add_resource_name(var.self);
statement(layout_for_variable(var), variable_decl(var), ";");
statement(layout_for_variable(var), to_qualifiers_glsl(var.self),
variable_decl(type, to_name(var.self), var.self), ";");
}
}
}
@ -2486,7 +2487,7 @@ string CompilerGLSL::enclose_expression(const string &expr)
if (!expr.empty())
{
auto c = expr.front();
if (c == '-' || c == '+' || c == '!' || c == '~')
if (c == '-' || c == '+' || c == '!' || c == '~' || c == '&' || c == '*')
need_parens = true;
}
@ -2520,6 +2521,28 @@ string CompilerGLSL::enclose_expression(const string &expr)
return expr;
}
string CompilerGLSL::dereference_expression(const std::string &expr)
{
// If this expression starts with an address-of operator ('&'), then
// just return the part after the operator.
// TODO: Strip parens if unnecessary?
if (expr.at(0) == '&')
return expr.substr(1);
else
return join('*', expr);
}
string CompilerGLSL::address_of_expression(const std::string &expr)
{
// If this expression starts with a dereference operator ('*'), then
// just return the part after the operator.
// TODO: Strip parens if unnecessary?
if (expr.at(0) == '*')
return expr.substr(1);
else
return join('&', expr);
}
// Just like to_expression except that we enclose the expression inside parentheses if needed.
string CompilerGLSL::to_enclosed_expression(uint32_t id, bool register_expression_read)
{
@ -2548,6 +2571,33 @@ string CompilerGLSL::to_enclosed_unpacked_expression(uint32_t id)
return to_enclosed_expression(id);
}
string CompilerGLSL::to_dereferenced_expression(uint32_t id, bool register_expression_read)
{
auto &type = expression_type(id);
if (type.pointer && should_dereference(id))
return dereference_expression(to_enclosed_expression(id, register_expression_read));
else
return to_expression(id, register_expression_read);
}
string CompilerGLSL::to_pointer_expression(uint32_t id)
{
auto &type = expression_type(id);
if (type.pointer && expression_is_lvalue(id) && !should_dereference(id))
return address_of_expression(to_enclosed_expression(id));
else
return to_expression(id);
}
string CompilerGLSL::to_enclosed_pointer_expression(uint32_t id)
{
auto &type = expression_type(id);
if (type.pointer && expression_is_lvalue(id) && !should_dereference(id))
return address_of_expression(to_enclosed_expression(id));
else
return to_enclosed_expression(id);
}
string CompilerGLSL::to_extract_component_expression(uint32_t id, uint32_t index)
{
auto expr = to_enclosed_expression(id);
@ -2860,10 +2910,14 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop)
string CompilerGLSL::constant_expression(const SPIRConstant &c)
{
if (!c.subconstants.empty())
{
auto &type = get<SPIRType>(c.constant_type);
auto &type = get<SPIRType>(c.constant_type);
if (type.pointer)
{
return backend.null_pointer_literal;
}
else if (!c.subconstants.empty())
{
// Handles Arrays and structures.
string res;
if (backend.use_initializer_list && backend.use_typed_initializer_list && type.basetype == SPIRType::Struct &&
@ -3909,8 +3963,8 @@ string CompilerGLSL::to_ternary_expression(const SPIRType &restype, uint32_t sel
auto &lerptype = expression_type(select);
if (lerptype.vecsize == 1)
expr = join(to_enclosed_expression(select), " ? ", to_enclosed_expression(true_value), " : ",
to_enclosed_expression(false_value));
expr = join(to_enclosed_expression(select), " ? ", to_enclosed_pointer_expression(true_value), " : ",
to_enclosed_pointer_expression(false_value));
else
{
auto swiz = [this](uint32_t expression, uint32_t i) { return to_extract_component_expression(expression, i); };
@ -3938,6 +3992,13 @@ void CompilerGLSL::emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left,
auto &lerptype = expression_type(lerp);
auto &restype = get<SPIRType>(result_type);
// If this results in a variable pointer, assume it may be written through.
if (restype.pointer)
{
register_write(left);
register_write(right);
}
string mix_op;
bool has_boolean_mix = backend.boolean_mix_support &&
((options.es && options.version >= 310) || (!options.es && options.version >= 450));
@ -5484,8 +5545,8 @@ const char *CompilerGLSL::index_to_swizzle(uint32_t index)
}
string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count,
bool index_is_literal, bool chain_only, AccessChainMeta *meta,
bool register_expression_read)
bool index_is_literal, bool chain_only, bool ptr_chain,
AccessChainMeta *meta, bool register_expression_read)
{
string expr;
if (!chain_only)
@ -5507,8 +5568,59 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
{
uint32_t index = indices[i];
const auto append_index = [&]() {
expr += "[";
if (index_is_literal)
expr += convert_to_string(index);
else
expr += to_expression(index, register_expression_read);
expr += "]";
};
// Pointer chains
if (ptr_chain && i == 0)
{
// If we are flattening multidimensional arrays, only create opening bracket on first
// array index.
if (options.flatten_multidimensional_arrays)
{
dimension_flatten = type->array.size() >= 1;
pending_array_enclose = dimension_flatten;
if (pending_array_enclose)
expr += "[";
}
if (options.flatten_multidimensional_arrays && dimension_flatten)
{
// If we are flattening multidimensional arrays, do manual stride computation.
if (index_is_literal)
expr += convert_to_string(index);
else
expr += to_enclosed_expression(index, register_expression_read);
for (auto j = uint32_t(type->array.size()); j; j--)
{
expr += " * ";
expr += enclose_expression(to_array_size(*type, j - 1));
}
if (type->array.empty())
pending_array_enclose = false;
else
expr += " + ";
if (!pending_array_enclose)
expr += "]";
}
else
{
append_index();
}
access_chain_is_arrayed = true;
}
// Arrays
if (!type->array.empty())
else if (!type->array.empty())
{
// If we are flattening multidimensional arrays, only create opening bracket on first
// array index.
@ -5522,15 +5634,6 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
assert(type->parent_type);
const auto append_index = [&]() {
expr += "[";
if (index_is_literal)
expr += convert_to_string(index);
else
expr += to_expression(index, register_expression_read);
expr += "]";
};
auto *var = maybe_get<SPIRVariable>(base);
if (backend.force_gl_in_out_block && i == 0 && var && is_builtin_variable(*var) &&
!has_decoration(type->self, DecorationBlock))
@ -5625,7 +5728,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice
if (!qual_mbr_name.empty())
expr = qual_mbr_name;
else
expr += to_member_reference(maybe_get_backing_variable(base), *type, index);
expr += to_member_reference(base, *type, index, ptr_chain);
}
if (has_member_decoration(type->self, index, DecorationInvariant))
@ -5713,13 +5816,14 @@ string CompilerGLSL::to_flattened_struct_member(const SPIRVariable &var, uint32_
}
string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type,
AccessChainMeta *meta)
AccessChainMeta *meta, bool ptr_chain)
{
if (flattened_buffer_blocks.count(base))
{
uint32_t matrix_stride = 0;
bool need_transpose = false;
flattened_access_chain_offset(expression_type(base), indices, count, 0, 16, &need_transpose, &matrix_stride);
flattened_access_chain_offset(expression_type(base), indices, count, 0, 16, &need_transpose, &matrix_stride,
ptr_chain);
if (meta)
{
@ -5731,7 +5835,7 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32
}
else if (flattened_structs.count(base) && count > 0)
{
auto chain = access_chain_internal(base, indices, count, false, true, nullptr, false).substr(1);
auto chain = access_chain_internal(base, indices, count, false, true, ptr_chain, nullptr, false).substr(1);
if (meta)
{
meta->need_transpose = false;
@ -5741,7 +5845,7 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32
}
else
{
return access_chain_internal(base, indices, count, false, false, meta, false);
return access_chain_internal(base, indices, count, false, false, ptr_chain, meta, false);
}
}
@ -5935,11 +6039,9 @@ std::string CompilerGLSL::flattened_access_chain_vector(uint32_t base, const uin
}
}
std::pair<std::string, uint32_t> CompilerGLSL::flattened_access_chain_offset(const SPIRType &basetype,
const uint32_t *indices, uint32_t count,
uint32_t offset, uint32_t word_stride,
bool *need_transpose,
uint32_t *out_matrix_stride)
std::pair<std::string, uint32_t> CompilerGLSL::flattened_access_chain_offset(
const SPIRType &basetype, const uint32_t *indices, uint32_t count, uint32_t offset, uint32_t word_stride,
bool *need_transpose, uint32_t *out_matrix_stride, bool ptr_chain)
{
// Start traversing type hierarchy at the proper non-pointer types.
const auto *type = &get_non_pointer_type(basetype);
@ -5962,8 +6064,40 @@ std::pair<std::string, uint32_t> CompilerGLSL::flattened_access_chain_offset(con
{
uint32_t index = indices[i];
// Pointers
if (ptr_chain && i == 0)
{
// Here, the pointer type will be decorated with an array stride.
uint32_t array_stride = get_decoration(basetype.self, DecorationArrayStride);
if (!array_stride)
SPIRV_CROSS_THROW("SPIR-V does not define ArrayStride for buffer block.");
auto *constant = maybe_get<SPIRConstant>(index);
if (constant)
{
// Constant array access.
offset += constant->scalar() * array_stride;
}
else
{
// Dynamic array access.
if (array_stride % word_stride)
{
SPIRV_CROSS_THROW(
"Array stride for dynamic indexing must be divisible by the size of a 4-component vector. "
"Likely culprit here is a float or vec2 array inside a push constant block which is std430. "
"This cannot be flattened. Try using std140 layout instead.");
}
expr += to_enclosed_expression(index);
expr += " * ";
expr += convert_to_string(array_stride / word_stride);
expr += " + ";
}
// Type ID is unchanged.
}
// Arrays
if (!type->array.empty())
else if (!type->array.empty())
{
// Here, the type_id will be a type ID for the array type itself.
uint32_t array_stride = get_decoration(type_id, DecorationArrayStride);
@ -6097,6 +6231,29 @@ std::pair<std::string, uint32_t> CompilerGLSL::flattened_access_chain_offset(con
return std::make_pair(expr, offset);
}
bool CompilerGLSL::should_dereference(uint32_t id)
{
const auto &type = expression_type(id);
// Non-pointer expressions don't need to be dereferenced.
if (!type.pointer)
return false;
// Handles shouldn't be dereferenced either.
if (!expression_is_lvalue(id))
return false;
// If id is a variable but not a phi variable, we should not dereference it.
if (auto *var = maybe_get<SPIRVariable>(id))
return var->phi_variable;
// If id is an access chain, we should not dereference it.
if (auto *expr = maybe_get<SPIRExpression>(id))
return !expr->access_chain;
// Otherwise, we should dereference this pointer expression.
return true;
}
bool CompilerGLSL::should_forward(uint32_t id)
{
// If id is a variable we will try to forward it regardless of force_temporary check below
@ -6560,7 +6717,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
// If we are forwarding this load,
// don't register the read to access chain here, defer that to when we actually use the expression,
// using the add_implied_read_expression mechanism.
auto expr = to_expression(ptr, !forward);
auto expr = to_dereferenced_expression(ptr, !forward);
// We might need to bitcast in order to load from a builtin.
bitcast_from_builtin_load(ptr, expr, get<SPIRType>(result_type));
@ -6591,6 +6748,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
case OpInBoundsAccessChain:
case OpAccessChain:
case OpPtrAccessChain:
{
auto *var = maybe_get<SPIRVariable>(ops[2]);
if (var)
@ -6599,13 +6757,15 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
// If the base is immutable, the access chain pointer must also be.
// If an expression is mutable and forwardable, we speculate that it is immutable.
AccessChainMeta meta;
auto e = access_chain(ops[2], &ops[3], length - 3, get<SPIRType>(ops[0]), &meta);
bool ptr_chain = opcode == OpPtrAccessChain;
auto e = access_chain(ops[2], &ops[3], length - 3, get<SPIRType>(ops[0]), &meta, ptr_chain);
auto &expr = set<SPIRExpression>(ops[1], move(e), ops[0], should_forward(ops[2]));
auto *backing_variable = maybe_get_backing_variable(ops[2]);
expr.loaded_from = backing_variable ? backing_variable->self : ops[2];
expr.need_transpose = meta.need_transpose;
expr.access_chain = true;
// Mark the result as being packed. Some platforms handled packed vectors differently than non-packed.
if (meta.storage_is_packed)
@ -6646,14 +6806,14 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
}
else
{
auto rhs = to_expression(ops[1]);
auto rhs = to_pointer_expression(ops[1]);
// Statements to OpStore may be empty if it is a struct with zero members. Just forward the store to /dev/null.
if (!rhs.empty())
{
handle_store_to_invariant_variable(ops[0], ops[1]);
auto lhs = to_expression(ops[0]);
auto lhs = to_dereferenced_expression(ops[0]);
// We might need to bitcast in order to store to a builtin.
bitcast_to_builtin_store(ops[0], rhs, expression_type(ops[1]));
@ -6667,6 +6827,10 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
register_write(ops[0]);
}
}
// Storing a pointer results in a variable pointer, so we must conservatively assume
// we can write through it.
if (expression_type(ops[1]).pointer)
register_write(ops[1]);
break;
}
@ -6961,14 +7125,14 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
//
// Including the base will prevent this and would trigger multiple reads
// from expression causing it to be forced to an actual temporary in GLSL.
auto expr = access_chain_internal(ops[2], &ops[3], length, true, true, &meta);
auto expr = access_chain_internal(ops[2], &ops[3], length, true, true, false, &meta);
e = &emit_op(result_type, id, expr, true, !expression_is_forwarded(ops[2]));
inherit_expression_dependencies(id, ops[2]);
e->base_expression = ops[2];
}
else
{
auto expr = access_chain_internal(ops[2], &ops[3], length, true, false, &meta);
auto expr = access_chain_internal(ops[2], &ops[3], length, true, false, false, &meta);
e = &emit_op(result_type, id, expr, should_forward(ops[2]), !expression_is_forwarded(ops[2]));
inherit_expression_dependencies(id, ops[2]);
}
@ -8623,7 +8787,7 @@ string CompilerGLSL::to_member_name(const SPIRType &type, uint32_t index)
return join("_m", index);
}
string CompilerGLSL::to_member_reference(const SPIRVariable *, const SPIRType &type, uint32_t index)
string CompilerGLSL::to_member_reference(uint32_t, const SPIRType &type, uint32_t index, bool)
{
return join(".", to_member_name(type, index));
}
@ -8853,7 +9017,7 @@ string CompilerGLSL::to_initializer_expression(const SPIRVariable &var)
string CompilerGLSL::variable_decl(const SPIRVariable &variable)
{
// Ignore the pointer type since GLSL doesn't have pointers.
auto &type = get<SPIRType>(variable.basetype);
auto &type = get_variable_data_type(variable);
if (type.pointer_depth > 1)
SPIRV_CROSS_THROW("Cannot declare pointer-to-pointer types.");
@ -9652,7 +9816,7 @@ void CompilerGLSL::flush_phi(uint32_t from, uint32_t to)
if (temporary_phi_variables.count(phi.local_variable))
rhs = join("_", phi.local_variable, "_copy");
else
rhs = to_expression(phi.local_variable);
rhs = to_pointer_expression(phi.local_variable);
if (!optimize_read_modify_write(get<SPIRType>(var.basetype), lhs, rhs))
statement(lhs, " = ", rhs, ";");
@ -9974,18 +10138,23 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block)
}
else
{
auto &var = get<SPIRVariable>(loop_var);
auto &type = get_variable_data_type(var);
if (expr.empty())
{
// For loop initializers are of the form <type id = value, id = value, id = value, etc ...
auto &var = get<SPIRVariable>(loop_var);
auto &type = get<SPIRType>(var.basetype);
expr = join(to_qualifiers_glsl(var.self), type_to_glsl(type), " ");
}
else
{
expr += ", ";
// In MSL, being based on C++, the asterisk marking a pointer
// binds to the identifier, not the type.
if (type.pointer)
expr += "* ";
}
auto &v = get<SPIRVariable>(loop_var);
expr += join(to_name(loop_var), " = ", to_expression(v.static_expression));
expr += join(to_name(loop_var), " = ", to_pointer_expression(var.static_expression));
}
}
return expr;

16
spirv_glsl.hpp
View File

@ -356,6 +356,7 @@ protected:
struct BackendVariations
{
std::string discard_literal = "discard";
std::string null_pointer_literal = "";
bool float_literal_suffix = false;
bool double_literal_suffix = true;
bool uint32_t_literal_suffix = true;
@ -417,6 +418,7 @@ protected:
void flush_variable_declaration(uint32_t id);
void flush_undeclared_variables(SPIRBlock &block);
bool should_dereference(uint32_t id);
bool should_forward(uint32_t id);
void emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left, uint32_t right, uint32_t lerp);
bool to_trivial_mix_op(const SPIRType &type, std::string &op, uint32_t left, uint32_t right, uint32_t lerp);
@ -445,10 +447,10 @@ protected:
SPIRExpression &emit_op(uint32_t result_type, uint32_t result_id, const std::string &rhs, bool forward_rhs,
bool suppress_usage_tracking = false);
std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, bool index_is_literal,
bool chain_only = false, AccessChainMeta *meta = nullptr,
bool chain_only = false, bool ptr_chain = false, AccessChainMeta *meta = nullptr,
bool register_expression_read = true);
std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type,
AccessChainMeta *meta = nullptr);
AccessChainMeta *meta = nullptr, bool ptr_chain = false);
std::string flattened_access_chain(uint32_t base, const uint32_t *indices, uint32_t count,
const SPIRType &target_type, uint32_t offset, uint32_t matrix_stride,
@ -464,7 +466,8 @@ protected:
std::pair<std::string, uint32_t> flattened_access_chain_offset(const SPIRType &basetype, const uint32_t *indices,
uint32_t count, uint32_t offset,
uint32_t word_stride, bool *need_transpose = nullptr,
uint32_t *matrix_stride = nullptr);
uint32_t *matrix_stride = nullptr,
bool ptr_chain = false);
const char *index_to_swizzle(uint32_t index);
std::string remap_swizzle(const SPIRType &result_type, uint32_t input_components, const std::string &expr);
@ -474,11 +477,16 @@ protected:
std::string to_enclosed_expression(uint32_t id, bool register_expression_read = true);
std::string to_unpacked_expression(uint32_t id);
std::string to_enclosed_unpacked_expression(uint32_t id);
std::string to_dereferenced_expression(uint32_t id, bool register_expression_read = true);
std::string to_pointer_expression(uint32_t id);
std::string to_enclosed_pointer_expression(uint32_t id);
std::string to_extract_component_expression(uint32_t id, uint32_t index);
std::string enclose_expression(const std::string &expr);
std::string dereference_expression(const std::string &expr);
std::string address_of_expression(const std::string &expr);
void strip_enclosed_expression(std::string &expr);
std::string to_member_name(const SPIRType &type, uint32_t index);
virtual std::string to_member_reference(const SPIRVariable *var, const SPIRType &type, uint32_t index);
virtual std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain);
std::string type_to_glsl_constructor(const SPIRType &type);
std::string argument_decl(const SPIRFunction::Parameter &arg);
virtual std::string to_qualifiers_glsl(uint32_t id);

123
spirv_msl.cpp
View File

@ -365,7 +365,7 @@ void CompilerMSL::emit_entry_point_declarations()
for (uint32_t array_id : buffer_arrays)
{
const auto &var = get<SPIRVariable>(array_id);
const auto &type = get<SPIRType>(var.basetype);
const auto &type = get_variable_data_type(var);
string name = get_name(array_id);
statement(get_argument_address_space(var) + " " + type_to_glsl(type) + "* " + name + "[] =");
begin_scope();
@ -387,6 +387,7 @@ string CompilerMSL::compile()
options.vulkan_semantics = true;
options.es = false;
options.version = 450;
backend.null_pointer_literal = "nullptr";
backend.float_literal_suffix = false;
backend.half_literal_suffix = "h";
backend.uint32_t_literal_suffix = true;
@ -604,6 +605,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
case OpLoad:
case OpInBoundsAccessChain:
case OpAccessChain:
case OpPtrAccessChain:
{
uint32_t base_id = ops[2];
if (global_var_ids.find(base_id) != global_var_ids.end())
@ -654,6 +656,16 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
break;
}
case OpSelect:
{
uint32_t base_id = ops[3];
if (global_var_ids.find(base_id) != global_var_ids.end())
added_arg_ids.insert(base_id);
base_id = ops[4];
if (global_var_ids.find(base_id) != global_var_ids.end())
added_arg_ids.insert(base_id);
}
default:
break;
}
@ -699,6 +711,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
ptr.self = mbr_type_id;
ptr.storage = var.storage;
ptr.pointer = true;
ptr.parent_type = mbr_type_id;
func.add_parameter(mbr_type_id, var_id, true);
set<SPIRVariable>(var_id, ptr_type_id, StorageClassFunction);
@ -863,7 +876,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage)
for (auto p_var : vars)
{
uint32_t type_id = p_var->basetype;
auto &type = get<SPIRType>(type_id);
auto &type = get_variable_data_type(*p_var);
if (type.basetype == SPIRType::Struct)
{
@ -1029,6 +1042,8 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage)
}
auto *usable_type = &type;
if (usable_type->pointer)
usable_type = &get<SPIRType>(usable_type->parent_type);
while (is_array(*usable_type) || is_matrix(*usable_type))
usable_type = &get<SPIRType>(usable_type->parent_type);
@ -1102,7 +1117,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage)
uint32_t ib_mbr_idx = uint32_t(ib_type.member_types.size());
type_id = ensure_correct_builtin_type(type_id, builtin);
p_var->basetype = type_id;
ib_type.member_types.push_back(type_id);
ib_type.member_types.push_back(get_non_pointer_type_id(type_id));
// Give the member a name
string mbr_name = ensure_valid_name(to_expression(p_var->self), "m");
@ -1120,7 +1135,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage)
{
type_id = ensure_correct_attribute_type(type_id, locn);
p_var->basetype = type_id;
ib_type.member_types[ib_mbr_idx] = type_id;
ib_type.member_types[ib_mbr_idx] = get_non_pointer_type_id(type_id);
}
set_member_decoration(ib_type_id, ib_mbr_idx, DecorationLocation, locn);
mark_location_as_used_by_shader(locn, storage);
@ -2778,7 +2793,7 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
string exp = string(op) + "(";
auto &type = expression_type(obj);
auto &type = get_non_pointer_type(expression_type(obj));
exp += "(volatile ";
auto *var = maybe_get_backing_variable(obj);
if (!var)
@ -3947,7 +3962,7 @@ string CompilerMSL::func_type_decl(SPIRType &type)
if (stage_out_var_id && ep_should_return_output)
{
auto &so_var = get<SPIRVariable>(stage_out_var_id);
auto &so_type = get<SPIRType>(so_var.basetype);
auto &so_type = get_variable_data_type(so_var);
return_type = type_to_glsl(so_type) + type_to_array_glsl(type);
}
@ -3975,7 +3990,7 @@ string CompilerMSL::func_type_decl(SPIRType &type)
return entry_type + " " + return_type;
}
// In MSL, address space qualifiers are required for all pointer or reference arguments
// In MSL, address space qualifiers are required for all pointer or reference variables
string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
{
const auto &type = get<SPIRType>(argument.basetype);
@ -4019,6 +4034,43 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
return "thread";
}
string CompilerMSL::get_type_address_space(const SPIRType &type)
{
switch (type.storage)
{
case StorageClassWorkgroup:
return "threadgroup";
case StorageClassStorageBuffer:
// FIXME: Need to use 'const device' for pointers into non-writable SSBOs
return "device";
case StorageClassUniform:
case StorageClassUniformConstant:
case StorageClassPushConstant:
if (type.basetype == SPIRType::Struct)
{
bool ssbo = has_decoration(type.self, DecorationBufferBlock);
// FIXME: Need to use 'const device' for pointers into non-writable SSBOs
if (ssbo)
return "device";
else
return "constant";
}
break;
case StorageClassFunction:
case StorageClassGeneric:
// No address space for plain values.
return type.pointer ? "thread" : "";
default:
break;
}
return "thread";
}
// Returns a string containing a comma-delimited list of args for the entry point function
string CompilerMSL::entry_point_args(bool append_comma)
{
@ -4028,7 +4080,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
if (stage_in_var_id)
{
auto &var = get<SPIRVariable>(stage_in_var_id);
auto &type = get<SPIRType>(var.basetype);
auto &type = get_variable_data_type(var);
if (!ep_args.empty())
ep_args += ", ";
@ -4054,7 +4106,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
if (id.get_type() == TypeVariable)
{
auto &var = id.get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
auto &type = get_variable_data_type(var);
uint32_t var_id = var.self;
@ -4090,7 +4142,7 @@ string CompilerMSL::entry_point_args(bool append_comma)
for (auto &r : resources)
{
auto &var = r.id->get<SPIRVariable>();
auto &type = get<SPIRType>(var.basetype);
auto &type = get_variable_data_type(var);
uint32_t var_id = var.self;
@ -4275,9 +4327,11 @@ 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);
auto &type = get_variable_data_type(var);
auto &var_type = get<SPIRType>(arg.type);
StorageClass storage = var_type.storage;
bool is_pointer = var_type.pointer;
// If we need to modify the name of the variable, make sure we use the original variable.
// Our alias is just a shadow variable.
@ -4285,7 +4339,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
if (arg.alias_global_variable && var.basevariable)
name_id = var.basevariable;
bool constref = !arg.alias_global_variable && type.pointer && arg.write_count == 0;
bool constref = !arg.alias_global_variable && is_pointer && arg.write_count == 0;
bool type_is_image = type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage ||
type.basetype == SPIRType::Sampler;
@ -4301,13 +4355,15 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
bool builtin = is_builtin_variable(var);
if (builtin)
decl += builtin_type_decl(static_cast<BuiltIn>(get_decoration(arg.id, DecorationBuiltIn)));
else if ((storage == StorageClassUniform || storage == StorageClassStorageBuffer) && is_array(type))
decl += join(type_to_glsl(type, arg.id), "*");
else
decl += type_to_glsl(type, arg.id);
bool opaque_handle = type.storage == StorageClassUniformConstant;
bool opaque_handle = storage == StorageClassUniformConstant;
if (!builtin && !opaque_handle && !type.pointer &&
(type.storage == StorageClassFunction || type.storage == StorageClassGeneric))
if (!builtin && !opaque_handle && !is_pointer &&
(storage == StorageClassFunction || storage == StorageClassGeneric))
{
// If the argument is a pure value and not an opaque type, we will pass by value.
decl += " ";
@ -4316,11 +4372,6 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
else if (is_array(type) && !type_is_image)
{
// Arrays of images and samplers are special cased.
if (get<SPIRVariable>(name_id).storage == StorageClassUniform ||
get<SPIRVariable>(name_id).storage == StorageClassStorageBuffer)
// If an array of buffers, declare an array of pointers, since we
// can't have an array of references.
decl += "*";
decl += " (&";
decl += to_expression(name_id);
decl += ")";
@ -4440,11 +4491,14 @@ void CompilerMSL::replace_illegal_names()
CompilerGLSL::replace_illegal_names();
}
string CompilerMSL::to_member_reference(const SPIRVariable *var, const SPIRType &type, uint32_t index)
string CompilerMSL::to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain)
{
auto *var = maybe_get<SPIRVariable>(base);
// If this is a buffer array, we have to dereference the buffer pointers.
if (var && (var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer) &&
!get<SPIRType>(var->basetype).array.empty())
// Otherwise, if this is a pointer expression, dereference it.
if ((var && ((var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer) &&
is_array(get<SPIRType>(var->basetype)))) ||
(!ptr_chain && should_dereference(base)))
return join("->", to_member_name(type, index));
else
return join(".", to_member_name(type, index));
@ -4466,10 +4520,27 @@ string CompilerMSL::to_qualifiers_glsl(uint32_t id)
// depend on a specific object's use of that type.
string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
{
// Ignore the pointer type since GLSL doesn't have pointers.
string type_name;
// Pointer?
if (type.pointer)
{
type_name = join(get_type_address_space(type), " ", type_to_glsl(get<SPIRType>(type.parent_type), id));
switch (type.basetype)
{
case SPIRType::Image:
case SPIRType::SampledImage:
case SPIRType::Sampler:
// These are handles.
break;
default:
// Anything else can be a raw pointer.
type_name += "*";
break;
}
return type_name;
}
switch (type.basetype)
{
case SPIRType::Struct:

3
spirv_msl.hpp
View File

@ -341,7 +341,7 @@ protected:
std::string unpack_expression_type(std::string expr_str, const SPIRType &type) override;
std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override;
bool skip_argument(uint32_t id) const override;
std::string to_member_reference(const SPIRVariable *var, const SPIRType &type, uint32_t index) override;
std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override;
std::string to_qualifiers_glsl(uint32_t id) override;
void replace_illegal_names() override;
void declare_undefined_values() override;
@ -391,6 +391,7 @@ protected:
bool is_member_packable(SPIRType &ib_type, uint32_t index);
MSLStructMemberKey get_struct_member_key(uint32_t type_id, uint32_t index);
std::string get_argument_address_space(const SPIRVariable &argument);
std::string get_type_address_space(const SPIRType &type);
void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1,
uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
bool op1_is_pointer = false, bool op1_is_literal = false, uint32_t op2 = 0);

7
spirv_parser.cpp
View File

@ -1081,7 +1081,12 @@ void Parser::make_constant_null(uint32_t id, uint32_t type)
{
auto &constant_type = get<SPIRType>(type);
if (!constant_type.array.empty())
if (constant_type.pointer)
{
auto &constant = set<SPIRConstant>(id, type);
constant.make_null(constant_type);
}
else if (!constant_type.array.empty())
{
assert(constant_type.parent_type);
uint32_t parent_id = ir.increase_bound_by(1);