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

Merge pull request #836 from KhronosGroup/fix-834

Browse Source 
MSL: Deal with resource name aliasing.
This commit is contained in:
Hans-Kristian Arntzen 2019-01-18 17:18:59 +01:00 committed by GitHub
commit 7430e78a8c
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
4 changed files with 89 additions and 0 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
reference/opt/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp Normal file
View File

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

23
reference/shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp Normal file
View File

@ -0,0 +1,23 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
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);
}

47
shaders-msl/asm/comp/struct-resource-name-aliasing.asm.comp Normal file
View File

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

3
spirv_msl.cpp
View File

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