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

Merge pull request #589 from KhronosGroup/fix-580

Browse Source 
Use const device address space for readonly SSBOs in MSL.
This commit is contained in:
Hans-Kristian Arntzen 2018-05-25 09:45:52 +01:00 committed by GitHub
commit db1ed375b0
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
27 changed files with 105 additions and 29 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
main.cpp
View File

@ -729,7 +729,8 @@ static int main_inner(int argc, char *argv[])
});
cbs.add("--remove-unused-variables", [&args](CLIParser &) { args.remove_unused = true; });
cbs.add("--combined-samplers-inherit-bindings", [&args](CLIParser &) { args.combined_samplers_inherit_bindings = true; });
cbs.add("--combined-samplers-inherit-bindings",
[&args](CLIParser &) { args.combined_samplers_inherit_bindings = true; });
cbs.default_handler = [&args](const char *value) { args.input = value; };
cbs.error_handler = [] { print_help(); };

2
reference/opt/shaders-msl/comp/basic.comp
View File

@ -21,7 +21,7 @@ struct SSBO3
uint counter;
};
kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
float4 _29 = _23.in_data[gl_GlobalInvocationID.x];
if (dot(_29, float4(1.0, 5.0, 6.0, 2.0)) > 8.19999980926513671875)

2
reference/opt/shaders-msl/comp/culling.comp
View File

@ -23,7 +23,7 @@ struct SSBO3
uint count;
};
kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
float _28 = _22.in_data[gl_GlobalInvocationID.x];
if (_28 > 12.0)

2
reference/opt/shaders-msl/comp/dowhile.comp
View File

@ -14,7 +14,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
int i = 0;
float4 _56;

2
reference/opt/shaders-msl/comp/inverse.comp
View File

@ -114,7 +114,7 @@ float2x2 spvInverse2x2(float2x2 m)
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
}
kernel void main0(device MatrixOut& _15 [[buffer(0)]], device MatrixIn& _20 [[buffer(1)]])
kernel void main0(device MatrixOut& _15 [[buffer(0)]], const device MatrixIn& _20 [[buffer(1)]])
{
_15.m2out = spvInverse2x2(_20.m2in);
_15.m3out = spvInverse3x3(_20.m3in);

2
reference/opt/shaders-msl/comp/mod.comp
View File

@ -22,7 +22,7 @@ Tx mod(Tx x, Ty y)
return x - y * floor(x / y);
}
kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
_33.out_data[gl_GlobalInvocationID.x] = mod(_23.in_data[gl_GlobalInvocationID.x], _33.out_data[gl_GlobalInvocationID.x]);
_33.out_data[gl_GlobalInvocationID.x] = as_type<float4>(as_type<uint4>(_23.in_data[gl_GlobalInvocationID.x]) % as_type<uint4>(_33.out_data[gl_GlobalInvocationID.x]));

2
reference/opt/shaders-msl/comp/modf.comp
View File

@ -13,7 +13,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
float4 i;
float4 _31 = modf(_23.in_data[gl_GlobalInvocationID.x], i);

2
reference/opt/shaders-msl/comp/read-write-only.comp
View File

@ -21,7 +21,7 @@ struct SSBO1
float4 data3;
};
kernel void main0(device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
kernel void main0(const device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
{
_10.data4 = _15.data0 + _21.data2;
_10.data5 = _15.data1 + _21.data3;

2
reference/opt/shaders-msl/comp/shared.comp
View File

@ -15,7 +15,7 @@ struct SSBO2
float out_data[1];
};
kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
{
threadgroup float sShared[4];
sShared[gl_LocalInvocationIndex] = _22.in_data[gl_GlobalInvocationID.x];

2
reference/opt/shaders-msl/comp/struct-layout.comp
View File

@ -18,7 +18,7 @@ struct SSBO
Foo in_data[1];
};
kernel void main0(device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
_23.out_data[gl_GlobalInvocationID.x].m = _30.in_data[gl_GlobalInvocationID.x].m * _30.in_data[gl_GlobalInvocationID.x].m;
}

2
reference/opt/shaders-msl/comp/torture-loop.comp
View File

@ -14,7 +14,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
float4 _99;
_99 = _24.in_data[gl_GlobalInvocationID.x];

22
reference/opt/shaders-msl/frag/readonly-ssbo.frag Normal file
View File

@ -0,0 +1,22 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO
{
float4 v;
};
struct main0_out
{
float4 FragColor [[color(0)]];
};
fragment main0_out main0(const device SSBO& _13 [[buffer(0)]])
{
main0_out out = {};
out.FragColor = _13.v + _13.v;
return out;
}

2
reference/shaders-msl-no-opt/comp/loop.comp
View File

@ -14,7 +14,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _177 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _177 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
float4 idat = _24.in_data[ident];

2
reference/shaders-msl/comp/basic.comp
View File

@ -21,7 +21,7 @@ struct SSBO3
uint counter;
};
kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _45 [[buffer(1)]], device SSBO3& _48 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
float4 idata = _23.in_data[ident];

2
reference/shaders-msl/comp/culling.comp
View File

@ -23,7 +23,7 @@ struct SSBO3
uint count;
};
kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _38 [[buffer(1)]], device SSBO3& _41 [[buffer(2)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
float idata = _22.in_data[ident];

2
reference/shaders-msl/comp/dowhile.comp
View File

@ -14,7 +14,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _28 [[buffer(0)]], device SSBO2& _52 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
int i = 0;

2
reference/shaders-msl/comp/inverse.comp
View File

@ -114,7 +114,7 @@ float2x2 spvInverse2x2(float2x2 m)
return (det != 0.0f) ? (adj * (1.0f / det)) : m;
}
kernel void main0(device MatrixOut& _15 [[buffer(0)]], device MatrixIn& _20 [[buffer(1)]])
kernel void main0(device MatrixOut& _15 [[buffer(0)]], const device MatrixIn& _20 [[buffer(1)]])
{
_15.m2out = spvInverse2x2(_20.m2in);
_15.m3out = spvInverse3x3(_20.m3in);

2
reference/shaders-msl/comp/mod.comp
View File

@ -22,7 +22,7 @@ Tx mod(Tx x, Ty y)
return x - y * floor(x / y);
}
kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _33 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
float4 v = mod(_23.in_data[ident], _33.out_data[ident]);

2
reference/shaders-msl/comp/modf.comp
View File

@ -13,7 +13,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _23 [[buffer(0)]], device SSBO2& _35 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
float4 i;

2
reference/shaders-msl/comp/read-write-only.comp
View File

@ -21,7 +21,7 @@ struct SSBO1
float4 data3;
};
kernel void main0(device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
kernel void main0(const device SSBO0& _15 [[buffer(0)]], device SSBO1& _21 [[buffer(1)]], device SSBO2& _10 [[buffer(2)]])
{
_10.data4 = _15.data0 + _21.data2;
_10.data5 = _15.data1 + _21.data3;

2
reference/shaders-msl/comp/shared.comp
View File

@ -15,7 +15,7 @@ struct SSBO2
float out_data[1];
};
kernel void main0(device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
kernel void main0(const device SSBO& _22 [[buffer(0)]], device SSBO2& _44 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]], uint gl_LocalInvocationIndex [[thread_index_in_threadgroup]])
{
threadgroup float sShared[4];
uint ident = gl_GlobalInvocationID.x;

2
reference/shaders-msl/comp/struct-layout.comp
View File

@ -18,7 +18,7 @@ struct SSBO
Foo in_data[1];
};
kernel void main0(device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _30 [[buffer(0)]], device SSBO2& _23 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
_23.out_data[ident].m = _30.in_data[ident].m * _30.in_data[ident].m;

2
reference/shaders-msl/comp/torture-loop.comp
View File

@ -14,7 +14,7 @@ struct SSBO2
float4 out_data[1];
};
kernel void main0(device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
kernel void main0(const device SSBO& _24 [[buffer(0)]], device SSBO2& _89 [[buffer(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]])
{
uint ident = gl_GlobalInvocationID.x;
float4 idat = _24.in_data[ident];

29
reference/shaders-msl/frag/readonly-ssbo.frag Normal file
View File

@ -0,0 +1,29 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct SSBO
{
float4 v;
};
struct main0_out
{
float4 FragColor [[color(0)]];
};
float4 read_from_function(const device SSBO& v_13)
{
return v_13.v;
}
fragment main0_out main0(const device SSBO& v_13 [[buffer(0)]])
{
main0_out out = {};
out.FragColor = v_13.v + read_from_function(v_13);
return out;
}

16
shaders-msl/frag/readonly-ssbo.frag Normal file
View File

@ -0,0 +1,16 @@
#version 450
layout(location = 0) out vec4 FragColor;
layout(binding = 0, std430) readonly buffer SSBO
{
vec4 v;
};
vec4 read_from_function()
{
return v;
}
void main()
{
FragColor = v + read_from_function();
}

2
spirv_cross_util.hpp
View File

@ -24,6 +24,6 @@ namespace spirv_cross_util
void rename_interface_variable(spirv_cross::Compiler &compiler, const std::vector<spirv_cross::Resource> &resources,
uint32_t location, const std::string &name);
void inherit_combined_sampler_bindings(spirv_cross::Compiler &compiler);
}
} // namespace spirv_cross_util
#endif

20
spirv_msl.cpp
View File

@ -3190,17 +3190,25 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument)
return "threadgroup";
case StorageClassStorageBuffer:
return "device";
{
auto flags = get_buffer_block_flags(argument);
return flags.get(DecorationNonWritable) ? "const device" : "device";
}
case StorageClassUniform:
case StorageClassUniformConstant:
case StorageClassPushConstant:
if (type.basetype == SPIRType::Struct)
return (meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock) &&
!meta[argument.self].decoration.decoration_flags.get(DecorationNonWritable)) ?
"device" :
"constant";
{
bool ssbo = has_decoration(type.self, DecorationBufferBlock);
if (!ssbo)
return "constant";
else
{
bool readonly = get_buffer_block_flags(argument).get(DecorationNonWritable);
return readonly ? "const device" : "device";
}
}
break;
default: