From 437fc87a893cb625bd7109b79b0a107adbb5fa07 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Fri, 18 Jan 2019 16:27:57 +0100 Subject: [PATCH] MSL: Deal with resource name aliasing. Apparently we didn't use those yet. MSL seems to be able to alias struct types and variable types to a degree, so that's why it has escaped testing until now. --- .../struct-resource-name-aliasing.asm.comp | 16 +++++++ .../struct-resource-name-aliasing.asm.comp | 23 +++++++++ .../struct-resource-name-aliasing.asm.comp | 47 +++++++++++++++++++ spirv_msl.cpp | 3 ++ 4 files changed, 89 insertions(+) create mode 100644 reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp create mode 100644 reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp create mode 100644 shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp diff --git a/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp b/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp new file mode 100644 index 00000000..5b1f0a03 --- /dev/null +++ b/reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp @@ -0,0 +1,16 @@ +#include +#include + +using namespace metal; + +struct bufA +{ + uint _data[1]; +}; + +kernel void main0(device bufA& bufA_1 [[buffer(0)]], device bufA& bufB [[buffer(1)]]) +{ + bufA_1._data[0] = 0u; + bufB._data[0] = 0u; +} + diff --git a/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp b/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp new file mode 100644 index 00000000..a3323bf2 --- /dev/null +++ b/reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp @@ -0,0 +1,23 @@ +#pragma clang diagnostic ignored "-Wmissing-prototypes" + +#include +#include + +using namespace metal; + +struct bufA +{ + uint _data[1]; +}; + +void _main(device bufA& bufA_1, device bufA& bufB) +{ + bufA_1._data[0] = 0u; + bufB._data[0] = 0u; +} + +kernel void main0(device bufA& bufA_1 [[buffer(0)]], device bufA& bufB [[buffer(1)]]) +{ + _main(bufA_1, bufB); +} + diff --git a/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp b/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp new file mode 100644 index 00000000..98d31537 --- /dev/null +++ b/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp @@ -0,0 +1,47 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 7 +; Bound: 21 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" + OpExecutionMode %main LocalSize 8 8 1 + OpSource HLSL 500 + OpName %main "main" + OpName %_main_ "@main(" + OpName %bufA "bufA" + OpMemberName %bufA 0 "@data" + OpName %bufA_0 "bufA" + OpName %bufB "bufB" + OpDecorate %_runtimearr_uint ArrayStride 4 + OpMemberDecorate %bufA 0 Offset 0 + OpDecorate %bufA BufferBlock + OpDecorate %bufA_0 DescriptorSet 0 + OpDecorate %bufB DescriptorSet 0 + %void = OpTypeVoid + %3 = OpTypeFunction %void + %uint = OpTypeInt 32 0 +%_runtimearr_uint = OpTypeRuntimeArray %uint + %bufA = OpTypeStruct %_runtimearr_uint +%_ptr_Uniform_bufA = OpTypePointer Uniform %bufA + %bufA_0 = OpVariable %_ptr_Uniform_bufA Uniform + %int = OpTypeInt 32 1 + %int_0 = OpConstant %int 0 + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %bufB = OpVariable %_ptr_Uniform_bufA Uniform + %main = OpFunction %void None %3 + %5 = OpLabel + %20 = OpFunctionCall %void %_main_ + OpReturn + OpFunctionEnd + %_main_ = OpFunction %void None %3 + %7 = OpLabel + %17 = OpAccessChain %_ptr_Uniform_uint %bufA_0 %int_0 %int_0 + OpStore %17 %uint_0 + %19 = OpAccessChain %_ptr_Uniform_uint %bufB %int_0 %int_0 + OpStore %19 %uint_0 + OpReturn + OpFunctionEnd diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 976722a8..98fe3ae0 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -4424,6 +4424,7 @@ string CompilerMSL::entry_point_args(bool append_comma) if (!ep_args.empty()) ep_args += ", "; + add_resource_name(var.self); ep_args += type_to_glsl(type) + " " + to_name(var.self) + " [[stage_in]]"; } @@ -4452,6 +4453,7 @@ string CompilerMSL::entry_point_args(bool append_comma) { if (type.basetype == SPIRType::SampledImage) { + add_resource_name(var_id); resources.push_back( { &id, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image) }); @@ -4464,6 +4466,7 @@ string CompilerMSL::entry_point_args(bool append_comma) else if (constexpr_samplers.count(var_id) == 0) { // constexpr samplers are not declared as resources. + add_resource_name(var_id); resources.push_back( { &id, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype) }); }