From 103817009ceb5952ea8cfb0a10d5955e3f921bf4 Mon Sep 17 00:00:00 2001 From: Chip Davis Date: Wed, 4 Sep 2019 13:57:17 -0500 Subject: [PATCH] MSL: Force storage images on iOS to use discrete descriptors. Writable textures cannot use argument buffers on iOS. They must be passed as arguments directly to the shader function. Since we won't know if a given storage image will have the `NonWritable` decoration at the time we encode the argument buffer, we must therefore pass all storage images as discrete arguments. Previously, we were throwing an error if we encountered an argument buffer with a writable texture in it on iOS. --- ...rs-image-load-store.ios.msl2.argument.comp | 11 +++++++++++ ...rs-image-load-store.ios.msl2.argument.comp | 11 +++++++++++ ...rs-image-load-store.ios.msl2.argument.comp | 10 ++++++++++ spirv_msl.cpp | 19 +++++++++++++------ spirv_msl.hpp | 6 +++--- 5 files changed, 48 insertions(+), 9 deletions(-) create mode 100644 reference/opt/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp create mode 100644 reference/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp create mode 100644 shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp diff --git a/reference/opt/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp b/reference/opt/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp new file mode 100644 index 00000000..25a0233a --- /dev/null +++ b/reference/opt/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp @@ -0,0 +1,11 @@ +#include +#include + +using namespace metal; + +kernel void main0(texture2d uImage [[texture(0)]], texture2d uImageRead [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + int2 _17 = int2(gl_GlobalInvocationID.xy); + uImage.write(uImageRead.read(uint2(_17)), uint2(_17)); +} + diff --git a/reference/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp b/reference/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp new file mode 100644 index 00000000..c23a9d1d --- /dev/null +++ b/reference/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp @@ -0,0 +1,11 @@ +#include +#include + +using namespace metal; + +kernel void main0(texture2d uImage [[texture(0)]], texture2d uImageRead [[texture(1)]], uint3 gl_GlobalInvocationID [[thread_position_in_grid]]) +{ + int2 coord = int2(gl_GlobalInvocationID.xy); + uImage.write(uImageRead.read(uint2(coord)), uint2(coord)); +} + diff --git a/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp b/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp new file mode 100644 index 00000000..72ca8899 --- /dev/null +++ b/shaders-msl/comp/argument-buffers-image-load-store.ios.msl2.argument.comp @@ -0,0 +1,10 @@ +#version 450 + +layout(set = 0, binding = 1, r32f) writeonly uniform image2D uImage; +layout(set = 0, binding = 2, r32f) readonly uniform image2D uImageRead; + +void main() +{ + ivec2 coord = ivec2(gl_GlobalInvocationID.xy); + imageStore(uImage, coord, imageLoad(uImageRead, coord)); +} diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 88199522..f0170258 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -8161,15 +8161,19 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) SmallVector resources; - ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + ir.for_each_typed_id([&](uint32_t var_id, SPIRVariable &var) { if ((var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) && !is_hidden_variable(var)) { auto &type = get_variable_data_type(var); - uint32_t var_id = var.self; - if (var.storage != StorageClassPushConstant) + // Very specifically, image load-store in argument buffers are disallowed on MSL on iOS. + // But we won't know when the argument buffer is encoded whether this image will have + // a NonWritable decoration. So just use discrete arguments for all storage images + // on iOS. + if (!(msl_options.is_ios() && type.basetype == SPIRType::Image && type.image.sampled == 2) && + var.storage != StorageClassPushConstant) { uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); if (descriptor_set_is_argument_buffer(desc_set)) @@ -11034,9 +11038,12 @@ void CompilerMSL::analyze_argument_buffers() else if (!constexpr_sampler) { // constexpr samplers are not declared as resources. - add_resource_name(var_id); - resources_in_set[desc_set].push_back( - { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype), 0 }); + if (!msl_options.is_ios() || type.basetype != SPIRType::Image || type.image.sampled != 2) + { + add_resource_name(var_id); + resources_in_set[desc_set].push_back( + { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype), 0 }); + } } // Check if this descriptor set needs a swizzle buffer. diff --git a/spirv_msl.hpp b/spirv_msl.hpp index b18b9b8c..fec4d34c 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -54,9 +54,9 @@ struct MSLVertexAttr // Matches the binding index of a MSL resource for a binding within a descriptor set. // Taken together, the stage, desc_set and binding combine to form a reference to a resource // descriptor used in a particular shading stage. -// If using MSL 2.0 argument buffers, and the descriptor set is not marked as a discrete descriptor set, -// the binding reference we remap to will become an [[id(N)]] attribute within -// the "descriptor set" argument buffer structure. +// If using MSL 2.0 argument buffers, the descriptor set is not marked as a discrete descriptor set, +// and (for iOS only) the resource is not a storage image (sampled != 2), the binding reference we +// remap to will become an [[id(N)]] attribute within the "descriptor set" argument buffer structure. // For resources which are bound in the "classic" MSL 1.0 way or discrete descriptors, the remap will become a // [[buffer(N)]], [[texture(N)]] or [[sampler(N)]] depending on the resource types used. struct MSLResourceBinding