|
|
@@ -201,13 +201,17 @@ bool CompilerMSL::is_var_runtime_size_array(const SPIRVariable &var) const
|
|
|
uint32_t CompilerMSL::get_resource_array_size(const SPIRType &type, uint32_t id) const
|
|
|
{
|
|
|
uint32_t array_size = to_array_size_literal(type);
|
|
|
- if (array_size)
|
|
|
+
|
|
|
+ // If we have argument buffers, we need to honor the ABI by using the correct array size
|
|
|
+ // from the layout. Only use shader declared size if we're not using argument buffers.
|
|
|
+ uint32_t desc_set = get_decoration(id, DecorationDescriptorSet);
|
|
|
+ if (!descriptor_set_is_argument_buffer(desc_set) && array_size)
|
|
|
return array_size;
|
|
|
|
|
|
- StageSetBinding tuple = { get_entry_point().model, get_decoration(id, DecorationDescriptorSet),
|
|
|
+ StageSetBinding tuple = { get_entry_point().model, desc_set,
|
|
|
get_decoration(id, DecorationBinding) };
|
|
|
auto itr = resource_bindings.find(tuple);
|
|
|
- return itr != end(resource_bindings) ? itr->second.first.count : 0;
|
|
|
+ return itr != end(resource_bindings) ? itr->second.first.count : array_size;
|
|
|
}
|
|
|
|
|
|
uint32_t CompilerMSL::get_automatic_msl_resource_binding(uint32_t id) const
|
|
|
@@ -267,11 +271,14 @@ void CompilerMSL::build_implicit_builtins()
|
|
|
active_input_builtins.get(BuiltInInstanceIndex) || active_input_builtins.get(BuiltInBaseInstance));
|
|
|
bool need_local_invocation_index = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInSubgroupId);
|
|
|
bool need_workgroup_size = msl_options.emulate_subgroups && active_input_builtins.get(BuiltInNumSubgroups);
|
|
|
+ bool force_frag_depth_passthrough =
|
|
|
+ get_execution_model() == ExecutionModelFragment && !uses_explicit_early_fragment_test() && need_subpass_input &&
|
|
|
+ msl_options.enable_frag_depth_builtin && msl_options.input_attachment_is_ds_attachment;
|
|
|
|
|
|
if (need_subpass_input || need_sample_pos || need_subgroup_mask || need_vertex_params || need_tesc_params ||
|
|
|
need_tese_params || need_multiview || need_dispatch_base || need_vertex_base_params || need_grid_params ||
|
|
|
needs_sample_id || needs_subgroup_invocation_id || needs_subgroup_size || needs_helper_invocation ||
|
|
|
- has_additional_fixed_sample_mask() || need_local_invocation_index || need_workgroup_size)
|
|
|
+ has_additional_fixed_sample_mask() || need_local_invocation_index || need_workgroup_size || force_frag_depth_passthrough)
|
|
|
{
|
|
|
bool has_frag_coord = false;
|
|
|
bool has_sample_id = false;
|
|
|
@@ -288,6 +295,7 @@ void CompilerMSL::build_implicit_builtins()
|
|
|
bool has_helper_invocation = false;
|
|
|
bool has_local_invocation_index = false;
|
|
|
bool has_workgroup_size = false;
|
|
|
+ bool has_frag_depth = false;
|
|
|
uint32_t workgroup_id_type = 0;
|
|
|
|
|
|
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
@@ -308,6 +316,13 @@ void CompilerMSL::build_implicit_builtins()
|
|
|
mark_implicit_builtin(StorageClassOutput, BuiltInSampleMask, var.self);
|
|
|
does_shader_write_sample_mask = true;
|
|
|
}
|
|
|
+
|
|
|
+ if (force_frag_depth_passthrough && builtin == BuiltInFragDepth)
|
|
|
+ {
|
|
|
+ builtin_frag_depth_id = var.self;
|
|
|
+ mark_implicit_builtin(StorageClassOutput, BuiltInFragDepth, var.self);
|
|
|
+ has_frag_depth = true;
|
|
|
+ }
|
|
|
}
|
|
|
|
|
|
if (var.storage != StorageClassInput)
|
|
|
@@ -898,6 +913,36 @@ void CompilerMSL::build_implicit_builtins()
|
|
|
builtin_workgroup_size_id = var_id;
|
|
|
mark_implicit_builtin(StorageClassInput, BuiltInWorkgroupSize, var_id);
|
|
|
}
|
|
|
+
|
|
|
+ if (!has_frag_depth && force_frag_depth_passthrough)
|
|
|
+ {
|
|
|
+ uint32_t offset = ir.increase_bound_by(3);
|
|
|
+ uint32_t type_id = offset;
|
|
|
+ uint32_t type_ptr_id = offset + 1;
|
|
|
+ uint32_t var_id = offset + 2;
|
|
|
+
|
|
|
+ // Create gl_FragDepth
|
|
|
+ SPIRType float_type { OpTypeFloat };
|
|
|
+ float_type.basetype = SPIRType::Float;
|
|
|
+ float_type.width = 32;
|
|
|
+ float_type.vecsize = 1;
|
|
|
+ set<SPIRType>(type_id, float_type);
|
|
|
+
|
|
|
+ SPIRType float_type_ptr_in = float_type;
|
|
|
+ float_type_ptr_in.op = spv::OpTypePointer;
|
|
|
+ float_type_ptr_in.pointer = true;
|
|
|
+ float_type_ptr_in.pointer_depth++;
|
|
|
+ float_type_ptr_in.parent_type = type_id;
|
|
|
+ float_type_ptr_in.storage = StorageClassOutput;
|
|
|
+
|
|
|
+ auto &ptr_in_type = set<SPIRType>(type_ptr_id, float_type_ptr_in);
|
|
|
+ ptr_in_type.self = type_id;
|
|
|
+ set<SPIRVariable>(var_id, type_ptr_id, StorageClassOutput);
|
|
|
+ set_decoration(var_id, DecorationBuiltIn, BuiltInFragDepth);
|
|
|
+ builtin_frag_depth_id = var_id;
|
|
|
+ mark_implicit_builtin(StorageClassOutput, BuiltInFragDepth, var_id);
|
|
|
+ active_output_builtins.set(BuiltInFragDepth);
|
|
|
+ }
|
|
|
}
|
|
|
|
|
|
if (needs_swizzle_buffer_def)
|
|
|
@@ -1314,48 +1359,29 @@ void CompilerMSL::emit_entry_point_declarations()
|
|
|
uint32_t arg_id = argument_buffer_ids[desc_set];
|
|
|
uint32_t base_index = dynamic_buffer.second.first;
|
|
|
|
|
|
- if (!type.array.empty())
|
|
|
+ if (is_array(type))
|
|
|
{
|
|
|
- // This is complicated, because we need to support arrays of arrays.
|
|
|
- // And it's even worse if the outermost dimension is a runtime array, because now
|
|
|
- // all this complicated goop has to go into the shader itself. (FIXME)
|
|
|
if (!type.array[type.array.size() - 1])
|
|
|
SPIRV_CROSS_THROW("Runtime arrays with dynamic offsets are not supported yet.");
|
|
|
- else
|
|
|
- {
|
|
|
- is_using_builtin_array = true;
|
|
|
- statement(get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name,
|
|
|
- type_to_array_glsl(type), " =");
|
|
|
|
|
|
- uint32_t dim = uint32_t(type.array.size());
|
|
|
- uint32_t j = 0;
|
|
|
- for (SmallVector<uint32_t> indices(type.array.size());
|
|
|
- indices[type.array.size() - 1] < to_array_size_literal(type); j++)
|
|
|
- {
|
|
|
- while (dim > 0)
|
|
|
- {
|
|
|
- begin_scope();
|
|
|
- --dim;
|
|
|
- }
|
|
|
+ is_using_builtin_array = true;
|
|
|
+ statement(get_argument_address_space(var), " ", type_to_glsl(type), "* ", to_restrict(var_id, true), name,
|
|
|
+ type_to_array_glsl(type, var_id), " =");
|
|
|
|
|
|
- string arrays;
|
|
|
- for (uint32_t i = uint32_t(type.array.size()); i; --i)
|
|
|
- arrays += join("[", indices[i - 1], "]");
|
|
|
- statement("(", get_argument_address_space(var), " ", type_to_glsl(type), "* ",
|
|
|
- to_restrict(var_id, false), ")((", get_argument_address_space(var), " char* ",
|
|
|
- to_restrict(var_id, false), ")", to_name(arg_id), ".", ensure_valid_name(name, "m"),
|
|
|
- arrays, " + ", to_name(dynamic_offsets_buffer_id), "[", base_index + j, "]),");
|
|
|
+ uint32_t array_size = to_array_size_literal(type);
|
|
|
+ begin_scope();
|
|
|
|
|
|
- while (++indices[dim] >= to_array_size_literal(type, dim) && dim < type.array.size() - 1)
|
|
|
- {
|
|
|
- end_scope(",");
|
|
|
- indices[dim++] = 0;
|
|
|
- }
|
|
|
- }
|
|
|
- end_scope_decl();
|
|
|
- statement_no_indent("");
|
|
|
- is_using_builtin_array = false;
|
|
|
+ for (uint32_t i = 0; i < array_size; i++)
|
|
|
+ {
|
|
|
+ statement("(", get_argument_address_space(var), " ", type_to_glsl(type), "* ",
|
|
|
+ to_restrict(var_id, false), ")((", get_argument_address_space(var), " char* ",
|
|
|
+ to_restrict(var_id, false), ")", to_name(arg_id), ".", ensure_valid_name(name, "m"),
|
|
|
+ "[", i, "]", " + ", to_name(dynamic_offsets_buffer_id), "[", base_index + i, "]),");
|
|
|
}
|
|
|
+
|
|
|
+ end_scope_decl();
|
|
|
+ statement_no_indent("");
|
|
|
+ is_using_builtin_array = false;
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
@@ -1469,7 +1495,7 @@ void CompilerMSL::emit_entry_point_declarations()
|
|
|
is_using_builtin_array = true;
|
|
|
statement(desc_addr_space, " auto& ", to_restrict(var_id, true), to_name(var_id), " = (", addr_space, " ",
|
|
|
type_to_glsl(type), "* ", desc_addr_space, " (&)",
|
|
|
- type_to_array_glsl(type), ")", ir.meta[alias_id].decoration.qualified_alias, ";");
|
|
|
+ type_to_array_glsl(type, var_id), ")", ir.meta[alias_id].decoration.qualified_alias, ";");
|
|
|
is_using_builtin_array = false;
|
|
|
}
|
|
|
}
|
|
|
@@ -1553,8 +1579,10 @@ string CompilerMSL::compile()
|
|
|
if (needs_manual_helper_invocation_updates() &&
|
|
|
(active_input_builtins.get(BuiltInHelperInvocation) || needs_helper_invocation))
|
|
|
{
|
|
|
- string discard_expr =
|
|
|
- join(builtin_to_glsl(BuiltInHelperInvocation, StorageClassInput), " = true, discard_fragment()");
|
|
|
+ string builtin_helper_invocation = builtin_to_glsl(BuiltInHelperInvocation, StorageClassInput);
|
|
|
+ string discard_expr = join(builtin_helper_invocation, " = true, discard_fragment()");
|
|
|
+ if (msl_options.force_fragment_with_side_effects_execution)
|
|
|
+ discard_expr = join("!", builtin_helper_invocation, " ? (", discard_expr, ") : (void)0");
|
|
|
backend.discard_literal = discard_expr;
|
|
|
backend.demote_literal = discard_expr;
|
|
|
}
|
|
|
@@ -1584,6 +1612,8 @@ string CompilerMSL::compile()
|
|
|
add_active_interface_variable(builtin_dispatch_base_id);
|
|
|
if (builtin_sample_mask_id)
|
|
|
add_active_interface_variable(builtin_sample_mask_id);
|
|
|
+ if (builtin_frag_depth_id)
|
|
|
+ add_active_interface_variable(builtin_frag_depth_id);
|
|
|
|
|
|
// Create structs to hold input, output and uniform variables.
|
|
|
// Do output first to ensure out. is declared at top of entry function.
|
|
|
@@ -1703,14 +1733,15 @@ void CompilerMSL::preprocess_op_codes()
|
|
|
|
|
|
// Fragment shaders that both write to storage resources and discard fragments
|
|
|
// need checks on the writes, to work around Metal allowing these writes despite
|
|
|
- // the fragment being dead.
|
|
|
- if (msl_options.check_discarded_frag_stores && preproc.uses_discard &&
|
|
|
- (preproc.uses_buffer_write || preproc.uses_image_write))
|
|
|
+ // the fragment being dead. We also require to force Metal to execute fragment
|
|
|
+ // shaders instead of being prematurely discarded.
|
|
|
+ if (preproc.uses_discard && (preproc.uses_buffer_write || preproc.uses_image_write))
|
|
|
{
|
|
|
- frag_shader_needs_discard_checks = true;
|
|
|
- needs_helper_invocation = true;
|
|
|
+ bool should_enable = (msl_options.check_discarded_frag_stores || msl_options.force_fragment_with_side_effects_execution);
|
|
|
+ frag_shader_needs_discard_checks |= msl_options.check_discarded_frag_stores;
|
|
|
+ needs_helper_invocation |= should_enable;
|
|
|
// Fragment discard store checks imply manual HelperInvocation updates.
|
|
|
- msl_options.manual_helper_invocation_updates = true;
|
|
|
+ msl_options.manual_helper_invocation_updates |= should_enable;
|
|
|
}
|
|
|
|
|
|
if (is_intersection_query())
|
|
|
@@ -1881,8 +1912,13 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
|
|
|
{
|
|
|
uint32_t base_id = ops[0];
|
|
|
if (global_var_ids.find(base_id) != global_var_ids.end())
|
|
|
+ {
|
|
|
added_arg_ids.insert(base_id);
|
|
|
|
|
|
+ if (msl_options.input_attachment_is_ds_attachment && base_id == builtin_frag_depth_id)
|
|
|
+ writes_to_depth = true;
|
|
|
+ }
|
|
|
+
|
|
|
uint32_t rvalue_id = ops[1];
|
|
|
if (global_var_ids.find(rvalue_id) != global_var_ids.end())
|
|
|
added_arg_ids.insert(rvalue_id);
|
|
|
@@ -2921,20 +2957,35 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
|
|
|
uint32_t mbr_idx, InterfaceBlockMeta &meta,
|
|
|
const string &mbr_name_qual,
|
|
|
const string &var_chain_qual,
|
|
|
- uint32_t &location, uint32_t &var_mbr_idx)
|
|
|
+ uint32_t &location, uint32_t &var_mbr_idx,
|
|
|
+ const Bitset &interpolation_qual)
|
|
|
{
|
|
|
auto &entry_func = get<SPIRFunction>(ir.default_entry_point);
|
|
|
|
|
|
BuiltIn builtin = BuiltInMax;
|
|
|
bool is_builtin = is_member_builtin(var_type, mbr_idx, &builtin);
|
|
|
- bool is_flat =
|
|
|
- has_member_decoration(var_type.self, mbr_idx, DecorationFlat) || has_decoration(var.self, DecorationFlat);
|
|
|
- bool is_noperspective = has_member_decoration(var_type.self, mbr_idx, DecorationNoPerspective) ||
|
|
|
+ bool is_flat = interpolation_qual.get(DecorationFlat) ||
|
|
|
+ has_member_decoration(var_type.self, mbr_idx, DecorationFlat) ||
|
|
|
+ has_decoration(var.self, DecorationFlat);
|
|
|
+ bool is_noperspective = interpolation_qual.get(DecorationNoPerspective) ||
|
|
|
+ has_member_decoration(var_type.self, mbr_idx, DecorationNoPerspective) ||
|
|
|
has_decoration(var.self, DecorationNoPerspective);
|
|
|
- bool is_centroid = has_member_decoration(var_type.self, mbr_idx, DecorationCentroid) ||
|
|
|
+ bool is_centroid = interpolation_qual.get(DecorationCentroid) ||
|
|
|
+ has_member_decoration(var_type.self, mbr_idx, DecorationCentroid) ||
|
|
|
has_decoration(var.self, DecorationCentroid);
|
|
|
- bool is_sample =
|
|
|
- has_member_decoration(var_type.self, mbr_idx, DecorationSample) || has_decoration(var.self, DecorationSample);
|
|
|
+ bool is_sample = interpolation_qual.get(DecorationSample) ||
|
|
|
+ has_member_decoration(var_type.self, mbr_idx, DecorationSample) ||
|
|
|
+ has_decoration(var.self, DecorationSample);
|
|
|
+
|
|
|
+ Bitset inherited_qual;
|
|
|
+ if (is_flat)
|
|
|
+ inherited_qual.set(DecorationFlat);
|
|
|
+ if (is_noperspective)
|
|
|
+ inherited_qual.set(DecorationNoPerspective);
|
|
|
+ if (is_centroid)
|
|
|
+ inherited_qual.set(DecorationCentroid);
|
|
|
+ if (is_sample)
|
|
|
+ inherited_qual.set(DecorationSample);
|
|
|
|
|
|
uint32_t mbr_type_id = var_type.member_types[mbr_idx];
|
|
|
auto &mbr_type = get<SPIRType>(mbr_type_id);
|
|
|
@@ -2998,7 +3049,7 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass
|
|
|
add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type,
|
|
|
var, mbr_type, sub_mbr_idx,
|
|
|
meta, mbr_name, var_chain,
|
|
|
- location, var_mbr_idx);
|
|
|
+ location, var_mbr_idx, inherited_qual);
|
|
|
// FIXME: Recursive structs and tessellation breaks here.
|
|
|
var_mbr_idx++;
|
|
|
}
|
|
|
@@ -3485,7 +3536,7 @@ void CompilerMSL::emit_local_masked_variable(const SPIRVariable &masked_var, boo
|
|
|
get_entry_point().output_vertices;
|
|
|
statement("threadgroup ", type_to_glsl(type), " ",
|
|
|
"spvStorage", to_name(masked_var.self), "[", max_num_instances, "]",
|
|
|
- type_to_array_glsl(type), ";");
|
|
|
+ type_to_array_glsl(type, 0), ";");
|
|
|
|
|
|
// Assign a threadgroup slice to each PrimitiveID.
|
|
|
// We assume here that workgroup size is rounded to 32,
|
|
|
@@ -3684,7 +3735,7 @@ void CompilerMSL::add_variable_to_interface_block(StorageClass storage, const st
|
|
|
add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type,
|
|
|
var, var_type, mbr_idx, meta,
|
|
|
mbr_name_qual, var_chain_qual,
|
|
|
- location, var_mbr_idx);
|
|
|
+ location, var_mbr_idx, {});
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
@@ -5564,9 +5615,8 @@ void CompilerMSL::emit_custom_templates()
|
|
|
// otherwise they will cause problems when linked together in a single Metallib.
|
|
|
void CompilerMSL::emit_custom_functions()
|
|
|
{
|
|
|
- for (uint32_t i = kArrayCopyMultidimMax; i >= 2; i--)
|
|
|
- if (spv_function_implementations.count(static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + i)))
|
|
|
- spv_function_implementations.insert(static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + i - 1));
|
|
|
+ if (spv_function_implementations.count(SPVFuncImplArrayCopyMultidim))
|
|
|
+ spv_function_implementations.insert(SPVFuncImplArrayCopy);
|
|
|
|
|
|
if (spv_function_implementations.count(SPVFuncImplDynamicImageSampler))
|
|
|
{
|
|
|
@@ -5677,11 +5727,7 @@ void CompilerMSL::emit_custom_functions()
|
|
|
break;
|
|
|
|
|
|
case SPVFuncImplArrayCopy:
|
|
|
- case SPVFuncImplArrayOfArrayCopy2Dim:
|
|
|
- case SPVFuncImplArrayOfArrayCopy3Dim:
|
|
|
- case SPVFuncImplArrayOfArrayCopy4Dim:
|
|
|
- case SPVFuncImplArrayOfArrayCopy5Dim:
|
|
|
- case SPVFuncImplArrayOfArrayCopy6Dim:
|
|
|
+ case SPVFuncImplArrayCopyMultidim:
|
|
|
{
|
|
|
// Unfortunately we cannot template on the address space, so combinatorial explosion it is.
|
|
|
static const char *function_name_tags[] = {
|
|
|
@@ -5704,36 +5750,19 @@ void CompilerMSL::emit_custom_functions()
|
|
|
|
|
|
for (uint32_t variant = 0; variant < 12; variant++)
|
|
|
{
|
|
|
- uint8_t dimensions = spv_func - SPVFuncImplArrayCopyMultidimBase;
|
|
|
- string tmp = "template<typename T";
|
|
|
- for (uint8_t i = 0; i < dimensions; i++)
|
|
|
- {
|
|
|
- tmp += ", uint ";
|
|
|
- tmp += 'A' + i;
|
|
|
- }
|
|
|
- tmp += ">";
|
|
|
- statement(tmp);
|
|
|
-
|
|
|
- string array_arg;
|
|
|
- for (uint8_t i = 0; i < dimensions; i++)
|
|
|
- {
|
|
|
- array_arg += "[";
|
|
|
- array_arg += 'A' + i;
|
|
|
- array_arg += "]";
|
|
|
- }
|
|
|
-
|
|
|
- statement("inline void spvArrayCopy", function_name_tags[variant], dimensions, "(",
|
|
|
- dst_address_space[variant], " T (&dst)", array_arg, ", ", src_address_space[variant],
|
|
|
- " T (&src)", array_arg, ")");
|
|
|
-
|
|
|
+ bool is_multidim = spv_func == SPVFuncImplArrayCopyMultidim;
|
|
|
+ const char* dim = is_multidim ? "[N][M]" : "[N]";
|
|
|
+ statement("template<typename T, uint N", is_multidim ? ", uint M>" : ">");
|
|
|
+ statement("inline void spvArrayCopy", function_name_tags[variant], "(",
|
|
|
+ dst_address_space[variant], " T (&dst)", dim, ", ",
|
|
|
+ src_address_space[variant], " T (&src)", dim, ")");
|
|
|
begin_scope();
|
|
|
- statement("for (uint i = 0; i < A; i++)");
|
|
|
+ statement("for (uint i = 0; i < N; i++)");
|
|
|
begin_scope();
|
|
|
-
|
|
|
- if (dimensions == 1)
|
|
|
- statement("dst[i] = src[i];");
|
|
|
+ if (is_multidim)
|
|
|
+ statement("spvArrayCopy", function_name_tags[variant], "(dst[i], src[i]);");
|
|
|
else
|
|
|
- statement("spvArrayCopy", function_name_tags[variant], dimensions - 1, "(dst[i], src[i]);");
|
|
|
+ statement("dst[i] = src[i];");
|
|
|
end_scope();
|
|
|
end_scope();
|
|
|
statement("");
|
|
|
@@ -6234,6 +6263,57 @@ void CompilerMSL::emit_custom_functions()
|
|
|
statement("");
|
|
|
break;
|
|
|
|
|
|
+ case SPVFuncImplGatherConstOffsets:
|
|
|
+ statement("// Wrapper function that processes a texture gather with a constant offset array.");
|
|
|
+ statement("template<typename T, template<typename, access = access::sample, typename = void> class Tex, "
|
|
|
+ "typename Toff, typename... Tp>");
|
|
|
+ statement("inline vec<T, 4> spvGatherConstOffsets(const thread Tex<T>& t, sampler s, "
|
|
|
+ "Toff coffsets, component c, Tp... params) METAL_CONST_ARG(c)");
|
|
|
+ begin_scope();
|
|
|
+ statement("vec<T, 4> rslts[4];");
|
|
|
+ statement("for (uint i = 0; i < 4; i++)");
|
|
|
+ begin_scope();
|
|
|
+ statement("switch (c)");
|
|
|
+ begin_scope();
|
|
|
+ // Work around texture::gather() requiring its component parameter to be a constant expression
|
|
|
+ statement("case component::x:");
|
|
|
+ statement(" rslts[i] = t.gather(s, spvForward<Tp>(params)..., coffsets[i], component::x);");
|
|
|
+ statement(" break;");
|
|
|
+ statement("case component::y:");
|
|
|
+ statement(" rslts[i] = t.gather(s, spvForward<Tp>(params)..., coffsets[i], component::y);");
|
|
|
+ statement(" break;");
|
|
|
+ statement("case component::z:");
|
|
|
+ statement(" rslts[i] = t.gather(s, spvForward<Tp>(params)..., coffsets[i], component::z);");
|
|
|
+ statement(" break;");
|
|
|
+ statement("case component::w:");
|
|
|
+ statement(" rslts[i] = t.gather(s, spvForward<Tp>(params)..., coffsets[i], component::w);");
|
|
|
+ statement(" break;");
|
|
|
+ end_scope();
|
|
|
+ end_scope();
|
|
|
+ // Pull all values from the i0j0 component of each gather footprint
|
|
|
+ statement("return vec<T, 4>(rslts[0].w, rslts[1].w, rslts[2].w, rslts[3].w);");
|
|
|
+ end_scope();
|
|
|
+ statement("");
|
|
|
+ break;
|
|
|
+
|
|
|
+ case SPVFuncImplGatherCompareConstOffsets:
|
|
|
+ statement("// Wrapper function that processes a texture gather with a constant offset array.");
|
|
|
+ statement("template<typename T, template<typename, access = access::sample, typename = void> class Tex, "
|
|
|
+ "typename Toff, typename... Tp>");
|
|
|
+ statement("inline vec<T, 4> spvGatherCompareConstOffsets(const thread Tex<T>& t, sampler s, "
|
|
|
+ "Toff coffsets, Tp... params)");
|
|
|
+ begin_scope();
|
|
|
+ statement("vec<T, 4> rslts[4];");
|
|
|
+ statement("for (uint i = 0; i < 4; i++)");
|
|
|
+ begin_scope();
|
|
|
+ statement(" rslts[i] = t.gather_compare(s, spvForward<Tp>(params)..., coffsets[i]);");
|
|
|
+ end_scope();
|
|
|
+ // Pull all values from the i0j0 component of each gather footprint
|
|
|
+ statement("return vec<T, 4>(rslts[0].w, rslts[1].w, rslts[2].w, rslts[3].w);");
|
|
|
+ end_scope();
|
|
|
+ statement("");
|
|
|
+ break;
|
|
|
+
|
|
|
case SPVFuncImplSubgroupBroadcast:
|
|
|
// Metal doesn't allow broadcasting boolean values directly, but we can work around that by broadcasting
|
|
|
// them as integers.
|
|
|
@@ -7426,14 +7506,14 @@ void CompilerMSL::emit_custom_functions()
|
|
|
statement("template<typename T>");
|
|
|
statement("struct spvDescriptorArray");
|
|
|
begin_scope();
|
|
|
- statement("spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(ptr)");
|
|
|
+ statement("spvDescriptorArray(const device spvDescriptor<T>* ptr) : ptr(&ptr->value)");
|
|
|
begin_scope();
|
|
|
end_scope();
|
|
|
statement("const device T& operator [] (size_t i) const");
|
|
|
begin_scope();
|
|
|
- statement("return ptr[i].value;");
|
|
|
+ statement("return ptr[i];");
|
|
|
end_scope();
|
|
|
- statement("const device spvDescriptor<T>* ptr;");
|
|
|
+ statement("const device T* ptr;");
|
|
|
end_scope_decl();
|
|
|
statement("");
|
|
|
}
|
|
|
@@ -7498,6 +7578,17 @@ void CompilerMSL::emit_custom_functions()
|
|
|
statement("");
|
|
|
break;
|
|
|
|
|
|
+ case SPVFuncImplTextureCast:
|
|
|
+ statement("template <typename T, typename U>");
|
|
|
+ statement("T spvTextureCast(U img)");
|
|
|
+ begin_scope();
|
|
|
+ // MSL complains if you try to cast the texture itself, but casting the reference type is ... ok? *shrug*
|
|
|
+ // Gotta go what you gotta do I suppose.
|
|
|
+ statement("return reinterpret_cast<thread const T &>(img);");
|
|
|
+ end_scope();
|
|
|
+ statement("");
|
|
|
+ break;
|
|
|
+
|
|
|
default:
|
|
|
break;
|
|
|
}
|
|
|
@@ -9428,32 +9519,12 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
|
|
|
uint32_t op1 = ops[3];
|
|
|
auto &type = get<SPIRType>(result_type);
|
|
|
auto input_type = opcode == OpSMulExtended ? int_type : uint_type;
|
|
|
- auto &output_type = get_type(result_type);
|
|
|
string cast_op0, cast_op1;
|
|
|
|
|
|
- auto expected_type = binary_op_bitcast_helper(cast_op0, cast_op1, input_type, op0, op1, false);
|
|
|
-
|
|
|
+ binary_op_bitcast_helper(cast_op0, cast_op1, input_type, op0, op1, false);
|
|
|
emit_uninitialized_temporary_expression(result_type, result_id);
|
|
|
-
|
|
|
- string mullo_expr, mulhi_expr;
|
|
|
- mullo_expr = join(cast_op0, " * ", cast_op1);
|
|
|
- mulhi_expr = join("mulhi(", cast_op0, ", ", cast_op1, ")");
|
|
|
-
|
|
|
- auto &low_type = get_type(output_type.member_types[0]);
|
|
|
- auto &high_type = get_type(output_type.member_types[1]);
|
|
|
- if (low_type.basetype != input_type)
|
|
|
- {
|
|
|
- expected_type.basetype = input_type;
|
|
|
- mullo_expr = join(bitcast_glsl_op(low_type, expected_type), "(", mullo_expr, ")");
|
|
|
- }
|
|
|
- if (high_type.basetype != input_type)
|
|
|
- {
|
|
|
- expected_type.basetype = input_type;
|
|
|
- mulhi_expr = join(bitcast_glsl_op(high_type, expected_type), "(", mulhi_expr, ")");
|
|
|
- }
|
|
|
-
|
|
|
- statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", mullo_expr, ";");
|
|
|
- statement(to_expression(result_id), ".", to_member_name(type, 1), " = ", mulhi_expr, ";");
|
|
|
+ statement(to_expression(result_id), ".", to_member_name(type, 0), " = ", cast_op0, " * ", cast_op1, ";");
|
|
|
+ statement(to_expression(result_id), ".", to_member_name(type, 1), " = mulhi(", cast_op0, ", ", cast_op1, ");");
|
|
|
break;
|
|
|
}
|
|
|
|
|
|
@@ -10020,15 +10091,7 @@ bool CompilerMSL::emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rh
|
|
|
// we cannot easily detect this case ahead of time since it's
|
|
|
// context dependent. We might have to force a recompile here
|
|
|
// if this is the only use of array copies in our shader.
|
|
|
- if (type.array.size() > 1)
|
|
|
- {
|
|
|
- if (type.array.size() > kArrayCopyMultidimMax)
|
|
|
- SPIRV_CROSS_THROW("Cannot support this many dimensions for arrays of arrays.");
|
|
|
- auto func = static_cast<SPVFuncImpl>(SPVFuncImplArrayCopyMultidimBase + type.array.size());
|
|
|
- add_spv_func_and_recompile(func);
|
|
|
- }
|
|
|
- else
|
|
|
- add_spv_func_and_recompile(SPVFuncImplArrayCopy);
|
|
|
+ add_spv_func_and_recompile(type.array.size() > 1 ? SPVFuncImplArrayCopyMultidim : SPVFuncImplArrayCopy);
|
|
|
|
|
|
const char *tag = nullptr;
|
|
|
if (lhs_is_thread_storage && is_constant)
|
|
|
@@ -10060,13 +10123,13 @@ bool CompilerMSL::emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rh
|
|
|
|
|
|
// Pass internal array of spvUnsafeArray<> into wrapper functions
|
|
|
if (lhs_is_array_template && rhs_is_array_template && !msl_options.force_native_arrays)
|
|
|
- statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ".elements, ", to_expression(rhs_id), ".elements);");
|
|
|
+ statement("spvArrayCopy", tag, "(", lhs, ".elements, ", to_expression(rhs_id), ".elements);");
|
|
|
if (lhs_is_array_template && !msl_options.force_native_arrays)
|
|
|
- statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ".elements, ", to_expression(rhs_id), ");");
|
|
|
+ statement("spvArrayCopy", tag, "(", lhs, ".elements, ", to_expression(rhs_id), ");");
|
|
|
else if (rhs_is_array_template && !msl_options.force_native_arrays)
|
|
|
- statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ", ", to_expression(rhs_id), ".elements);");
|
|
|
+ statement("spvArrayCopy", tag, "(", lhs, ", ", to_expression(rhs_id), ".elements);");
|
|
|
else
|
|
|
- statement("spvArrayCopy", tag, type.array.size(), "(", lhs, ", ", to_expression(rhs_id), ");");
|
|
|
+ statement("spvArrayCopy", tag, "(", lhs, ", ", to_expression(rhs_id), ");");
|
|
|
}
|
|
|
|
|
|
return true;
|
|
|
@@ -10213,7 +10276,35 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id,
|
|
|
if (split_index != string::npos)
|
|
|
{
|
|
|
auto coord = obj_expression.substr(split_index + 1);
|
|
|
- exp += join(obj_expression.substr(0, split_index), ".", op, "(");
|
|
|
+ auto image_expr = obj_expression.substr(0, split_index);
|
|
|
+
|
|
|
+ // Handle problem cases with sign where we need signed min/max on a uint image for example.
|
|
|
+ // It seems to work to cast the texture type itself, even if it is probably wildly outside of spec,
|
|
|
+ // but SPIR-V requires this to work.
|
|
|
+ if ((opcode == OpAtomicUMax || opcode == OpAtomicUMin ||
|
|
|
+ opcode == OpAtomicSMax || opcode == OpAtomicSMin) &&
|
|
|
+ type.basetype != expected_type)
|
|
|
+ {
|
|
|
+ auto *backing_var = maybe_get_backing_variable(obj);
|
|
|
+ if (backing_var)
|
|
|
+ {
|
|
|
+ add_spv_func_and_recompile(SPVFuncImplTextureCast);
|
|
|
+
|
|
|
+ const auto *backing_type = &get<SPIRType>(backing_var->basetype);
|
|
|
+ while (backing_type->op != OpTypeImage)
|
|
|
+ backing_type = &get<SPIRType>(backing_type->parent_type);
|
|
|
+
|
|
|
+ auto img_type = *backing_type;
|
|
|
+ auto tmp_type = type;
|
|
|
+ tmp_type.basetype = expected_type;
|
|
|
+ img_type.image.type = ir.increase_bound_by(1);
|
|
|
+ set<SPIRType>(img_type.image.type, tmp_type);
|
|
|
+
|
|
|
+ image_expr = join("spvTextureCast<", type_to_glsl(img_type, obj), ">(", image_expr, ")");
|
|
|
+ }
|
|
|
+ }
|
|
|
+
|
|
|
+ exp += join(image_expr, ".", op, "(");
|
|
|
if (ptr_type.storage == StorageClassImage && res_type->image.arrayed)
|
|
|
{
|
|
|
switch (res_type->image.dim)
|
|
|
@@ -10401,19 +10492,54 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop,
|
|
|
|
|
|
op = get_remapped_glsl_op(op);
|
|
|
|
|
|
+ auto &restype = get<SPIRType>(result_type);
|
|
|
+
|
|
|
switch (op)
|
|
|
{
|
|
|
case GLSLstd450Sinh:
|
|
|
- emit_unary_func_op(result_type, id, args[0], "fast::sinh");
|
|
|
+ if (restype.basetype == SPIRType::Half)
|
|
|
+ {
|
|
|
+ // MSL does not have overload for half. Force-cast back to half.
|
|
|
+ auto expr = join("half(fast::sinh(", to_unpacked_expression(args[0]), "))");
|
|
|
+ emit_op(result_type, id, expr, should_forward(args[0]));
|
|
|
+ inherit_expression_dependencies(id, args[0]);
|
|
|
+ }
|
|
|
+ else
|
|
|
+ emit_unary_func_op(result_type, id, args[0], "fast::sinh");
|
|
|
break;
|
|
|
case GLSLstd450Cosh:
|
|
|
- emit_unary_func_op(result_type, id, args[0], "fast::cosh");
|
|
|
+ if (restype.basetype == SPIRType::Half)
|
|
|
+ {
|
|
|
+ // MSL does not have overload for half. Force-cast back to half.
|
|
|
+ auto expr = join("half(fast::cosh(", to_unpacked_expression(args[0]), "))");
|
|
|
+ emit_op(result_type, id, expr, should_forward(args[0]));
|
|
|
+ inherit_expression_dependencies(id, args[0]);
|
|
|
+ }
|
|
|
+ else
|
|
|
+ emit_unary_func_op(result_type, id, args[0], "fast::cosh");
|
|
|
break;
|
|
|
case GLSLstd450Tanh:
|
|
|
- emit_unary_func_op(result_type, id, args[0], "precise::tanh");
|
|
|
+ if (restype.basetype == SPIRType::Half)
|
|
|
+ {
|
|
|
+ // MSL does not have overload for half. Force-cast back to half.
|
|
|
+ auto expr = join("half(fast::tanh(", to_unpacked_expression(args[0]), "))");
|
|
|
+ emit_op(result_type, id, expr, should_forward(args[0]));
|
|
|
+ inherit_expression_dependencies(id, args[0]);
|
|
|
+ }
|
|
|
+ else
|
|
|
+ emit_unary_func_op(result_type, id, args[0], "precise::tanh");
|
|
|
break;
|
|
|
case GLSLstd450Atan2:
|
|
|
- emit_binary_func_op(result_type, id, args[0], args[1], "precise::atan2");
|
|
|
+ if (restype.basetype == SPIRType::Half)
|
|
|
+ {
|
|
|
+ // MSL does not have overload for half. Force-cast back to half.
|
|
|
+ auto expr = join("half(fast::atan2(", to_unpacked_expression(args[0]), ", ", to_unpacked_expression(args[1]), "))");
|
|
|
+ emit_op(result_type, id, expr, should_forward(args[0]) && should_forward(args[1]));
|
|
|
+ inherit_expression_dependencies(id, args[0]);
|
|
|
+ inherit_expression_dependencies(id, args[1]);
|
|
|
+ }
|
|
|
+ else
|
|
|
+ emit_binary_func_op(result_type, id, args[0], args[1], "precise::atan2");
|
|
|
break;
|
|
|
case GLSLstd450InverseSqrt:
|
|
|
emit_unary_func_op(result_type, id, args[0], "rsqrt");
|
|
|
@@ -10809,7 +10935,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &)
|
|
|
decl += "thread ";
|
|
|
decl += type_to_glsl(type);
|
|
|
decl += " (&spvReturnValue)";
|
|
|
- decl += type_to_array_glsl(type);
|
|
|
+ decl += type_to_array_glsl(type, 0);
|
|
|
if (!func.arguments.empty())
|
|
|
decl += ", ";
|
|
|
}
|
|
|
@@ -10933,8 +11059,7 @@ string CompilerMSL::to_function_name(const TextureFunctionNameArguments &args)
|
|
|
is_dynamic_img_sampler = has_extended_decoration(var->self, SPIRVCrossDecorationDynamicImageSampler);
|
|
|
}
|
|
|
|
|
|
- // Special-case gather. We have to alter the component being looked up
|
|
|
- // in the swizzle case.
|
|
|
+ // Special-case gather. We have to alter the component being looked up in the swizzle case.
|
|
|
if (msl_options.swizzle_texture_samples && args.base.is_gather && !is_dynamic_img_sampler &&
|
|
|
(!constexpr_sampler || !constexpr_sampler->ycbcr_conversion_enable))
|
|
|
{
|
|
|
@@ -10943,6 +11068,16 @@ string CompilerMSL::to_function_name(const TextureFunctionNameArguments &args)
|
|
|
return is_compare ? "spvGatherCompareSwizzle" : "spvGatherSwizzle";
|
|
|
}
|
|
|
|
|
|
+ // Special-case gather with an array of offsets. We have to lower into 4 separate gathers.
|
|
|
+ if (args.has_array_offsets && !is_dynamic_img_sampler &&
|
|
|
+ (!constexpr_sampler || !constexpr_sampler->ycbcr_conversion_enable))
|
|
|
+ {
|
|
|
+ bool is_compare = comparison_ids.count(img);
|
|
|
+ add_spv_func_and_recompile(is_compare ? SPVFuncImplGatherCompareConstOffsets : SPVFuncImplGatherConstOffsets);
|
|
|
+ add_spv_func_and_recompile(SPVFuncImplForwardArgs);
|
|
|
+ return is_compare ? "spvGatherCompareConstOffsets" : "spvGatherConstOffsets";
|
|
|
+ }
|
|
|
+
|
|
|
auto *combined = maybe_get<SPIRCombinedImageSampler>(img);
|
|
|
|
|
|
// Texture reference
|
|
|
@@ -11123,6 +11258,10 @@ string CompilerMSL::to_function_args(const TextureFunctionArguments &args, bool
|
|
|
farg_str += to_expression(combined ? combined->image : img);
|
|
|
}
|
|
|
|
|
|
+ // Gathers with constant offsets call a special function, so include the texture.
|
|
|
+ if (args.has_array_offsets)
|
|
|
+ farg_str += to_expression(img);
|
|
|
+
|
|
|
// Sampler reference
|
|
|
if (!args.base.is_fetch)
|
|
|
{
|
|
|
@@ -11139,11 +11278,17 @@ string CompilerMSL::to_function_args(const TextureFunctionArguments &args, bool
|
|
|
used_swizzle_buffer = true;
|
|
|
}
|
|
|
|
|
|
- // Swizzled gather puts the component before the other args, to allow template
|
|
|
- // deduction to work.
|
|
|
- if (args.component && msl_options.swizzle_texture_samples)
|
|
|
+ // Const offsets gather puts the const offsets before the other args.
|
|
|
+ if (args.has_array_offsets)
|
|
|
{
|
|
|
- forward = should_forward(args.component);
|
|
|
+ forward = forward && should_forward(args.offset);
|
|
|
+ farg_str += ", " + to_expression(args.offset);
|
|
|
+ }
|
|
|
+
|
|
|
+ // Const offsets gather or swizzled gather puts the component before the other args.
|
|
|
+ if (args.component && (args.has_array_offsets || msl_options.swizzle_texture_samples))
|
|
|
+ {
|
|
|
+ forward = forward && should_forward(args.component);
|
|
|
farg_str += ", " + to_component_argument(args.component);
|
|
|
}
|
|
|
}
|
|
|
@@ -11554,7 +11699,7 @@ string CompilerMSL::to_function_args(const TextureFunctionArguments &args, bool
|
|
|
// Add offsets
|
|
|
string offset_expr;
|
|
|
const SPIRType *offset_type = nullptr;
|
|
|
- if (args.offset && !args.base.is_fetch)
|
|
|
+ if (args.offset && !args.base.is_fetch && !args.has_array_offsets)
|
|
|
{
|
|
|
forward = forward && should_forward(args.offset);
|
|
|
offset_expr = to_expression(args.offset);
|
|
|
@@ -11593,7 +11738,7 @@ string CompilerMSL::to_function_args(const TextureFunctionArguments &args, bool
|
|
|
}
|
|
|
}
|
|
|
|
|
|
- if (args.component)
|
|
|
+ if (args.component && !args.has_array_offsets)
|
|
|
{
|
|
|
// If 2D has gather component, ensure it also has an offset arg
|
|
|
if (imgtype.image.dim == Dim2D && offset_expr.empty())
|
|
|
@@ -12273,7 +12418,7 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
|
|
|
variable_storage_requires_stage_io(StorageClassInput)));
|
|
|
if (is_ib_in_out && is_member_builtin(type, index, &builtin))
|
|
|
is_using_builtin_array = true;
|
|
|
- array_type = type_to_array_glsl(physical_type);
|
|
|
+ array_type = type_to_array_glsl(physical_type, orig_id);
|
|
|
}
|
|
|
|
|
|
if (orig_id)
|
|
|
@@ -12843,14 +12988,14 @@ uint32_t CompilerMSL::get_or_allocate_builtin_output_member_location(spv::BuiltI
|
|
|
string CompilerMSL::func_type_decl(SPIRType &type)
|
|
|
{
|
|
|
// The regular function return type. If not processing the entry point function, that's all we need
|
|
|
- string return_type = type_to_glsl(type) + type_to_array_glsl(type);
|
|
|
+ string return_type = type_to_glsl(type) + type_to_array_glsl(type, 0);
|
|
|
if (!processing_entry_point)
|
|
|
return return_type;
|
|
|
|
|
|
// If an outgoing interface block has been defined, and it should be returned, override the entry point return type
|
|
|
bool ep_should_return_output = !get_is_rasterization_disabled();
|
|
|
if (stage_out_var_id && ep_should_return_output)
|
|
|
- return_type = type_to_glsl(get_stage_out_struct_type()) + type_to_array_glsl(type);
|
|
|
+ return_type = type_to_glsl(get_stage_out_struct_type()) + type_to_array_glsl(type, 0);
|
|
|
|
|
|
// Prepend a entry type, based on the execution model
|
|
|
string entry_type;
|
|
|
@@ -14416,16 +14561,33 @@ void CompilerMSL::fix_up_shader_inputs_outputs()
|
|
|
}
|
|
|
}
|
|
|
else if (var.storage == StorageClassOutput && get_execution_model() == ExecutionModelFragment &&
|
|
|
- is_builtin_variable(var) && active_output_builtins.get(bi_type) &&
|
|
|
- bi_type == BuiltInSampleMask && has_additional_fixed_sample_mask())
|
|
|
+ is_builtin_variable(var) && active_output_builtins.get(bi_type))
|
|
|
{
|
|
|
- // If the additional fixed sample mask was set, we need to adjust the sample_mask
|
|
|
- // output to reflect that. If the shader outputs the sample_mask itself too, we need
|
|
|
- // to AND the two masks to get the final one.
|
|
|
- string op_str = does_shader_write_sample_mask ? " &= " : " = ";
|
|
|
- entry_func.fixup_hooks_out.push_back([=]() {
|
|
|
- statement(to_expression(builtin_sample_mask_id), op_str, additional_fixed_sample_mask_str(), ";");
|
|
|
- });
|
|
|
+ switch (bi_type)
|
|
|
+ {
|
|
|
+ case BuiltInSampleMask:
|
|
|
+ if (has_additional_fixed_sample_mask())
|
|
|
+ {
|
|
|
+ // If the additional fixed sample mask was set, we need to adjust the sample_mask
|
|
|
+ // output to reflect that. If the shader outputs the sample_mask itself too, we need
|
|
|
+ // to AND the two masks to get the final one.
|
|
|
+ string op_str = does_shader_write_sample_mask ? " &= " : " = ";
|
|
|
+ entry_func.fixup_hooks_out.push_back([=]() {
|
|
|
+ statement(to_expression(builtin_sample_mask_id), op_str, additional_fixed_sample_mask_str(), ";");
|
|
|
+ });
|
|
|
+ }
|
|
|
+ break;
|
|
|
+ case BuiltInFragDepth:
|
|
|
+ if (msl_options.input_attachment_is_ds_attachment && !writes_to_depth)
|
|
|
+ {
|
|
|
+ entry_func.fixup_hooks_out.push_back([=]() {
|
|
|
+ statement(to_expression(builtin_frag_depth_id), " = ", to_expression(builtin_frag_coord_id), ".z;");
|
|
|
+ });
|
|
|
+ }
|
|
|
+ break;
|
|
|
+ default:
|
|
|
+ break;
|
|
|
+ }
|
|
|
}
|
|
|
});
|
|
|
}
|
|
|
@@ -14717,7 +14879,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
|
|
|
}
|
|
|
decl += to_expression(name_id);
|
|
|
decl += ")";
|
|
|
- decl += type_to_array_glsl(type);
|
|
|
+ decl += type_to_array_glsl(type, name_id);
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
@@ -14768,7 +14930,7 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg)
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
- auto array_size_decl = type_to_array_glsl(type);
|
|
|
+ auto array_size_decl = type_to_array_glsl(type, name_id);
|
|
|
if (array_size_decl.empty())
|
|
|
decl += "& ";
|
|
|
else
|
|
|
@@ -15508,7 +15670,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
|
|
|
return type_to_glsl(type, id, false);
|
|
|
}
|
|
|
|
|
|
-string CompilerMSL::type_to_array_glsl(const SPIRType &type)
|
|
|
+string CompilerMSL::type_to_array_glsl(const SPIRType &type, uint32_t variable_id)
|
|
|
{
|
|
|
// Allow Metal to use the array<T> template to make arrays a value type
|
|
|
switch (type.basetype)
|
|
|
@@ -15516,11 +15678,20 @@ string CompilerMSL::type_to_array_glsl(const SPIRType &type)
|
|
|
case SPIRType::AtomicCounter:
|
|
|
case SPIRType::ControlPointArray:
|
|
|
case SPIRType::RayQuery:
|
|
|
- return CompilerGLSL::type_to_array_glsl(type);
|
|
|
+ return CompilerGLSL::type_to_array_glsl(type, variable_id);
|
|
|
|
|
|
default:
|
|
|
if (type_is_array_of_pointers(type) || using_builtin_array())
|
|
|
- return CompilerGLSL::type_to_array_glsl(type);
|
|
|
+ {
|
|
|
+ const SPIRVariable *var = variable_id ? &get<SPIRVariable>(variable_id) : nullptr;
|
|
|
+ if (var && (var->storage == StorageClassUniform || var->storage == StorageClassStorageBuffer) &&
|
|
|
+ is_array(get_variable_data_type(*var)))
|
|
|
+ {
|
|
|
+ return join("[", get_resource_array_size(type, variable_id), "]");
|
|
|
+ }
|
|
|
+ else
|
|
|
+ return CompilerGLSL::type_to_array_glsl(type, variable_id);
|
|
|
+ }
|
|
|
else
|
|
|
return "";
|
|
|
}
|
|
|
@@ -15686,8 +15857,8 @@ string CompilerMSL::image_type_glsl(const SPIRType &type, uint32_t id, bool memb
|
|
|
|
|
|
string img_type_name;
|
|
|
|
|
|
- // Bypass pointers because we need the real image struct
|
|
|
- auto &img_type = get<SPIRType>(type.self).image;
|
|
|
+ auto &img_type = type.image;
|
|
|
+
|
|
|
if (is_depth_image(type, id))
|
|
|
{
|
|
|
switch (img_type.dim)
|
|
|
@@ -17925,7 +18096,7 @@ void CompilerMSL::emit_argument_buffer_aliased_descriptor(const SPIRVariable &al
|
|
|
is_using_builtin_array = true;
|
|
|
|
|
|
bool needs_post_cast_deref = !is_array(data_type);
|
|
|
- string ref_type = needs_post_cast_deref ? "&" : join("(&)", type_to_array_glsl(var_type));
|
|
|
+ string ref_type = needs_post_cast_deref ? "&" : join("(&)", type_to_array_glsl(var_type, aliased_var.self));
|
|
|
|
|
|
if (is_var_runtime_size_array(aliased_var))
|
|
|
{
|
|
|
@@ -18214,9 +18385,9 @@ void CompilerMSL::analyze_argument_buffers()
|
|
|
|
|
|
// If needed, synthesize and add padding members.
|
|
|
// member_index and next_arg_buff_index are incremented when padding members are added.
|
|
|
- if (msl_options.pad_argument_buffer_resources)
|
|
|
+ if (msl_options.pad_argument_buffer_resources && resource.overlapping_var_id == 0)
|
|
|
{
|
|
|
- auto &rez_bind = get_argument_buffer_resource(desc_set, next_arg_buff_index);
|
|
|
+ auto rez_bind = get_argument_buffer_resource(desc_set, next_arg_buff_index);
|
|
|
while (resource.index > next_arg_buff_index)
|
|
|
{
|
|
|
switch (rez_bind.basetype)
|
|
|
@@ -18252,11 +18423,14 @@ void CompilerMSL::analyze_argument_buffers()
|
|
|
default:
|
|
|
break;
|
|
|
}
|
|
|
+
|
|
|
+ // After padding, retrieve the resource again. It will either be more padding, or the actual resource.
|
|
|
+ rez_bind = get_argument_buffer_resource(desc_set, next_arg_buff_index);
|
|
|
}
|
|
|
|
|
|
// Adjust the number of slots consumed by current member itself.
|
|
|
// Use the count value from the app, instead of the shader, in case the
|
|
|
- // shader is only accesing part, or even one element, of the array.
|
|
|
+ // shader is only accessing part, or even one element, of the array.
|
|
|
next_arg_buff_index += rez_bind.count;
|
|
|
}
|
|
|
|
|
|
@@ -18391,7 +18565,7 @@ void CompilerMSL::analyze_argument_buffers()
|
|
|
// that matches the resource index of the argument buffer index.
|
|
|
// This is a two-step lookup, first lookup the resource binding number from the argument buffer index,
|
|
|
// then lookup the resource binding using the binding number.
|
|
|
-MSLResourceBinding &CompilerMSL::get_argument_buffer_resource(uint32_t desc_set, uint32_t arg_idx)
|
|
|
+const MSLResourceBinding &CompilerMSL::get_argument_buffer_resource(uint32_t desc_set, uint32_t arg_idx) const
|
|
|
{
|
|
|
auto stage = get_entry_point().model;
|
|
|
StageSetBinding arg_idx_tuple = { stage, desc_set, arg_idx };
|
Бранимир Караџић