|
|
@@ -1,5 +1,6 @@
|
|
|
/*
|
|
|
* Copyright 2016-2021 Robert Konrad
|
|
|
+ * SPDX-License-Identifier: Apache-2.0 OR MIT
|
|
|
*
|
|
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
|
|
* you may not use this file except in compliance with the License.
|
|
|
@@ -19,7 +20,6 @@
|
|
|
* At your option, you may choose to accept this material under either:
|
|
|
* 1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
|
|
|
* 2. The MIT License, found at <http://opensource.org/licenses/MIT>.
|
|
|
- * SPDX-License-Identifier: Apache-2.0 OR MIT.
|
|
|
*/
|
|
|
|
|
|
#include "spirv_hlsl.hpp"
|
|
|
@@ -570,7 +570,7 @@ void CompilerHLSL::emit_builtin_outputs_in_struct()
|
|
|
switch (builtin)
|
|
|
{
|
|
|
case BuiltInPosition:
|
|
|
- type = "float4";
|
|
|
+ type = is_position_invariant() && backend.support_precise_qualifier ? "precise float4" : "float4";
|
|
|
semantic = legacy ? "POSITION" : "SV_Position";
|
|
|
break;
|
|
|
|
|
|
@@ -641,7 +641,6 @@ void CompilerHLSL::emit_builtin_outputs_in_struct()
|
|
|
|
|
|
default:
|
|
|
SPIRV_CROSS_THROW("Unsupported builtin in HLSL.");
|
|
|
- break;
|
|
|
}
|
|
|
|
|
|
if (type && semantic)
|
|
|
@@ -770,7 +769,6 @@ void CompilerHLSL::emit_builtin_inputs_in_struct()
|
|
|
|
|
|
default:
|
|
|
SPIRV_CROSS_THROW("Unsupported builtin in HLSL.");
|
|
|
- break;
|
|
|
}
|
|
|
|
|
|
if (type && semantic)
|
|
|
@@ -818,8 +816,8 @@ string CompilerHLSL::to_interpolation_qualifiers(const Bitset &flags)
|
|
|
res += "patch "; // Seems to be different in actual HLSL.
|
|
|
if (flags.get(DecorationSample))
|
|
|
res += "sample ";
|
|
|
- if (flags.get(DecorationInvariant))
|
|
|
- res += "invariant "; // Not supported?
|
|
|
+ if (flags.get(DecorationInvariant) && backend.support_precise_qualifier)
|
|
|
+ res += "precise "; // Not supported?
|
|
|
|
|
|
return res;
|
|
|
}
|
|
|
@@ -854,48 +852,25 @@ std::string CompilerHLSL::to_initializer_expression(const SPIRVariable &var)
|
|
|
return CompilerGLSL::to_initializer_expression(var);
|
|
|
}
|
|
|
|
|
|
-void CompilerHLSL::emit_io_block(const SPIRVariable &var)
|
|
|
+void CompilerHLSL::emit_interface_block_member_in_struct(const SPIRVariable &var, uint32_t member_index,
|
|
|
+ uint32_t location,
|
|
|
+ std::unordered_set<uint32_t> &active_locations)
|
|
|
{
|
|
|
auto &execution = get_entry_point();
|
|
|
-
|
|
|
- auto &type = get<SPIRType>(var.basetype);
|
|
|
- add_resource_name(type.self);
|
|
|
-
|
|
|
- statement("struct ", to_name(type.self));
|
|
|
- begin_scope();
|
|
|
- type.member_name_cache.clear();
|
|
|
-
|
|
|
- uint32_t base_location = get_decoration(var.self, DecorationLocation);
|
|
|
-
|
|
|
- for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
|
|
|
- {
|
|
|
- string semantic;
|
|
|
- if (has_member_decoration(type.self, i, DecorationLocation))
|
|
|
- {
|
|
|
- uint32_t location = get_member_decoration(type.self, i, DecorationLocation);
|
|
|
- semantic = join(" : ", to_semantic(location, execution.model, var.storage));
|
|
|
- }
|
|
|
- else
|
|
|
- {
|
|
|
- // If the block itself has a location, but not its members, use the implicit location.
|
|
|
- // There could be a conflict if the block members partially specialize the locations.
|
|
|
- // It is unclear how SPIR-V deals with this. Assume this does not happen for now.
|
|
|
- uint32_t location = base_location + i;
|
|
|
- semantic = join(" : ", to_semantic(location, execution.model, var.storage));
|
|
|
- }
|
|
|
-
|
|
|
- add_member_name(type, i);
|
|
|
-
|
|
|
- auto &membertype = get<SPIRType>(type.member_types[i]);
|
|
|
- statement(to_interpolation_qualifiers(get_member_decoration_bitset(type.self, i)),
|
|
|
- variable_decl(membertype, to_member_name(type, i)), semantic, ";");
|
|
|
- }
|
|
|
-
|
|
|
- end_scope_decl();
|
|
|
- statement("");
|
|
|
-
|
|
|
- statement("static ", variable_decl(var), ";");
|
|
|
- statement("");
|
|
|
+ auto type = get<SPIRType>(var.basetype);
|
|
|
+ auto semantic = to_semantic(location, execution.model, var.storage);
|
|
|
+ auto mbr_name = join(to_name(type.self), "_", to_member_name(type, member_index));
|
|
|
+ auto &mbr_type = get<SPIRType>(type.member_types[member_index]);
|
|
|
+
|
|
|
+ statement(to_interpolation_qualifiers(get_member_decoration_bitset(type.self, member_index)),
|
|
|
+ type_to_glsl(mbr_type),
|
|
|
+ " ", mbr_name, type_to_array_glsl(mbr_type),
|
|
|
+ " : ", semantic, ";");
|
|
|
+
|
|
|
+ // Structs and arrays should consume more locations.
|
|
|
+ uint32_t consumed_locations = type_to_consumed_locations(mbr_type);
|
|
|
+ for (uint32_t i = 0; i < consumed_locations; i++)
|
|
|
+ active_locations.insert(location + i);
|
|
|
}
|
|
|
|
|
|
void CompilerHLSL::emit_interface_block_in_struct(const SPIRVariable &var, unordered_set<uint32_t> &active_locations)
|
|
|
@@ -930,7 +905,6 @@ void CompilerHLSL::emit_interface_block_in_struct(const SPIRVariable &var, unord
|
|
|
|
|
|
bool need_matrix_unroll = var.storage == StorageClassInput && execution.model == ExecutionModelVertex;
|
|
|
|
|
|
- auto &m = ir.meta[var.self].decoration;
|
|
|
auto name = to_name(var.self);
|
|
|
if (use_location_number)
|
|
|
{
|
|
|
@@ -938,8 +912,8 @@ void CompilerHLSL::emit_interface_block_in_struct(const SPIRVariable &var, unord
|
|
|
|
|
|
// If an explicit location exists, use it with TEXCOORD[N] semantic.
|
|
|
// Otherwise, pick a vacant location.
|
|
|
- if (m.decoration_flags.get(DecorationLocation))
|
|
|
- location_number = m.location;
|
|
|
+ if (has_decoration(var.self, DecorationLocation))
|
|
|
+ location_number = get_decoration(var.self, DecorationLocation);
|
|
|
else
|
|
|
location_number = get_vacant_location();
|
|
|
|
|
|
@@ -1188,10 +1162,10 @@ void CompilerHLSL::emit_composite_constants()
|
|
|
|
|
|
auto &type = this->get<SPIRType>(c.constant_type);
|
|
|
|
|
|
- // Cannot declare block type constants here.
|
|
|
- // We do not have the struct type yet.
|
|
|
- bool is_block = has_decoration(type.self, DecorationBlock);
|
|
|
- if (!is_block && (type.basetype == SPIRType::Struct || !type.array.empty()))
|
|
|
+ if (type.basetype == SPIRType::Struct && is_builtin_type(type))
|
|
|
+ return;
|
|
|
+
|
|
|
+ if (type.basetype == SPIRType::Struct || !type.array.empty())
|
|
|
{
|
|
|
auto name = to_name(c.self);
|
|
|
statement("static const ", variable_decl(type, name), " = ", constant_expression(c), ";");
|
|
|
@@ -1209,6 +1183,18 @@ void CompilerHLSL::emit_specialization_constants_and_structs()
|
|
|
SpecializationConstant wg_x, wg_y, wg_z;
|
|
|
ID workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z);
|
|
|
|
|
|
+ std::unordered_set<TypeID> io_block_types;
|
|
|
+ ir.for_each_typed_id<SPIRVariable>([&](uint32_t, const SPIRVariable &var) {
|
|
|
+ auto &type = this->get<SPIRType>(var.basetype);
|
|
|
+ if ((var.storage == StorageClassInput || var.storage == StorageClassOutput) &&
|
|
|
+ !var.remapped_variable && type.pointer && !is_builtin_variable(var) &&
|
|
|
+ interface_variable_exists_in_entry_point(var.self) &&
|
|
|
+ has_decoration(type.self, DecorationBlock))
|
|
|
+ {
|
|
|
+ io_block_types.insert(type.self);
|
|
|
+ }
|
|
|
+ });
|
|
|
+
|
|
|
auto loop_lock = ir.create_loop_hard_lock();
|
|
|
for (auto &id_ : ir.ids_for_constant_or_type)
|
|
|
{
|
|
|
@@ -1251,9 +1237,11 @@ void CompilerHLSL::emit_specialization_constants_and_structs()
|
|
|
else if (id.get_type() == TypeType)
|
|
|
{
|
|
|
auto &type = id.get<SPIRType>();
|
|
|
- if (type.basetype == SPIRType::Struct && type.array.empty() && !type.pointer &&
|
|
|
- (!ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) &&
|
|
|
- !ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock)))
|
|
|
+ bool is_non_io_block = has_decoration(type.self, DecorationBlock) &&
|
|
|
+ io_block_types.count(type.self) == 0;
|
|
|
+ bool is_buffer_block = has_decoration(type.self, DecorationBufferBlock);
|
|
|
+ if (type.basetype == SPIRType::Struct && type.array.empty() &&
|
|
|
+ !type.pointer && !is_non_io_block && !is_buffer_block)
|
|
|
{
|
|
|
if (emitted)
|
|
|
statement("");
|
|
|
@@ -1379,16 +1367,12 @@ void CompilerHLSL::emit_resources()
|
|
|
|
|
|
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
auto &type = this->get<SPIRType>(var.basetype);
|
|
|
- bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
|
|
|
|
|
|
- // Do not emit I/O blocks here.
|
|
|
- // I/O blocks can be arrayed, so we must deal with them separately to support geometry shaders
|
|
|
- // and tessellation down the line.
|
|
|
- if (!block && var.storage != StorageClassFunction && !var.remapped_variable && type.pointer &&
|
|
|
+ if (var.storage != StorageClassFunction && !var.remapped_variable && type.pointer &&
|
|
|
(var.storage == StorageClassInput || var.storage == StorageClassOutput) && !is_builtin_variable(var) &&
|
|
|
interface_variable_exists_in_entry_point(var.self))
|
|
|
{
|
|
|
- // Only emit non-builtins which are not blocks here. Builtin variables are handled separately.
|
|
|
+ // Builtin variables are handled separately.
|
|
|
emit_interface_block_globally(var);
|
|
|
emitted = true;
|
|
|
}
|
|
|
@@ -1402,69 +1386,72 @@ void CompilerHLSL::emit_resources()
|
|
|
require_output = false;
|
|
|
unordered_set<uint32_t> active_inputs;
|
|
|
unordered_set<uint32_t> active_outputs;
|
|
|
- SmallVector<SPIRVariable *> input_variables;
|
|
|
- SmallVector<SPIRVariable *> output_variables;
|
|
|
+
|
|
|
+ struct IOVariable
|
|
|
+ {
|
|
|
+ const SPIRVariable *var;
|
|
|
+ uint32_t location;
|
|
|
+ uint32_t block_member_index;
|
|
|
+ bool block;
|
|
|
+ };
|
|
|
+
|
|
|
+ SmallVector<IOVariable> input_variables;
|
|
|
+ SmallVector<IOVariable> output_variables;
|
|
|
+
|
|
|
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
auto &type = this->get<SPIRType>(var.basetype);
|
|
|
- bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
|
|
|
+ bool block = has_decoration(type.self, DecorationBlock);
|
|
|
|
|
|
if (var.storage != StorageClassInput && var.storage != StorageClassOutput)
|
|
|
return;
|
|
|
|
|
|
- // Do not emit I/O blocks here.
|
|
|
- // I/O blocks can be arrayed, so we must deal with them separately to support geometry shaders
|
|
|
- // and tessellation down the line.
|
|
|
- if (!block && !var.remapped_variable && type.pointer && !is_builtin_variable(var) &&
|
|
|
+ if (!var.remapped_variable && type.pointer && !is_builtin_variable(var) &&
|
|
|
interface_variable_exists_in_entry_point(var.self))
|
|
|
{
|
|
|
- if (var.storage == StorageClassInput)
|
|
|
- input_variables.push_back(&var);
|
|
|
- else
|
|
|
- output_variables.push_back(&var);
|
|
|
- }
|
|
|
-
|
|
|
- // Reserve input and output locations for block variables as necessary.
|
|
|
- if (block && !is_builtin_variable(var) && interface_variable_exists_in_entry_point(var.self))
|
|
|
- {
|
|
|
- auto &active = var.storage == StorageClassInput ? active_inputs : active_outputs;
|
|
|
- for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
|
|
|
+ if (block)
|
|
|
{
|
|
|
- if (has_member_decoration(type.self, i, DecorationLocation))
|
|
|
+ for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
|
|
|
{
|
|
|
- uint32_t location = get_member_decoration(type.self, i, DecorationLocation);
|
|
|
- active.insert(location);
|
|
|
+ uint32_t location = get_declared_member_location(var, i, false);
|
|
|
+ if (var.storage == StorageClassInput)
|
|
|
+ input_variables.push_back({ &var, location, i, true });
|
|
|
+ else
|
|
|
+ output_variables.push_back({ &var, location, i, true });
|
|
|
}
|
|
|
}
|
|
|
-
|
|
|
- // Emit the block struct and a global variable here.
|
|
|
- emit_io_block(var);
|
|
|
+ else
|
|
|
+ {
|
|
|
+ uint32_t location = get_decoration(var.self, DecorationLocation);
|
|
|
+ if (var.storage == StorageClassInput)
|
|
|
+ input_variables.push_back({ &var, location, 0, false });
|
|
|
+ else
|
|
|
+ output_variables.push_back({ &var, location, 0, false });
|
|
|
+ }
|
|
|
}
|
|
|
});
|
|
|
|
|
|
- const auto variable_compare = [&](const SPIRVariable *a, const SPIRVariable *b) -> bool {
|
|
|
+ const auto variable_compare = [&](const IOVariable &a, const IOVariable &b) -> bool {
|
|
|
// Sort input and output variables based on, from more robust to less robust:
|
|
|
// - Location
|
|
|
// - Variable has a location
|
|
|
// - Name comparison
|
|
|
// - Variable has a name
|
|
|
// - Fallback: ID
|
|
|
- bool has_location_a = has_decoration(a->self, DecorationLocation);
|
|
|
- bool has_location_b = has_decoration(b->self, DecorationLocation);
|
|
|
+ bool has_location_a = a.block || has_decoration(a.var->self, DecorationLocation);
|
|
|
+ bool has_location_b = b.block || has_decoration(b.var->self, DecorationLocation);
|
|
|
|
|
|
if (has_location_a && has_location_b)
|
|
|
- {
|
|
|
- return get_decoration(a->self, DecorationLocation) < get_decoration(b->self, DecorationLocation);
|
|
|
- }
|
|
|
+ return a.location < b.location;
|
|
|
else if (has_location_a && !has_location_b)
|
|
|
return true;
|
|
|
else if (!has_location_a && has_location_b)
|
|
|
return false;
|
|
|
|
|
|
- const auto &name1 = to_name(a->self);
|
|
|
- const auto &name2 = to_name(b->self);
|
|
|
+ const auto &name1 = to_name(a.var->self);
|
|
|
+ const auto &name2 = to_name(b.var->self);
|
|
|
|
|
|
if (name1.empty() && name2.empty())
|
|
|
- return a->self < b->self;
|
|
|
+ return a.var->self < b.var->self;
|
|
|
else if (name1.empty())
|
|
|
return true;
|
|
|
else if (name2.empty())
|
|
|
@@ -1491,8 +1478,13 @@ void CompilerHLSL::emit_resources()
|
|
|
|
|
|
begin_scope();
|
|
|
sort(input_variables.begin(), input_variables.end(), variable_compare);
|
|
|
- for (auto var : input_variables)
|
|
|
- emit_interface_block_in_struct(*var, active_inputs);
|
|
|
+ for (auto &var : input_variables)
|
|
|
+ {
|
|
|
+ if (var.block)
|
|
|
+ emit_interface_block_member_in_struct(*var.var, var.block_member_index, var.location, active_inputs);
|
|
|
+ else
|
|
|
+ emit_interface_block_in_struct(*var.var, active_inputs);
|
|
|
+ }
|
|
|
emit_builtin_inputs_in_struct();
|
|
|
end_scope_decl();
|
|
|
statement("");
|
|
|
@@ -1504,10 +1496,14 @@ void CompilerHLSL::emit_resources()
|
|
|
statement("struct SPIRV_Cross_Output");
|
|
|
|
|
|
begin_scope();
|
|
|
- // FIXME: Use locations properly if they exist.
|
|
|
sort(output_variables.begin(), output_variables.end(), variable_compare);
|
|
|
- for (auto var : output_variables)
|
|
|
- emit_interface_block_in_struct(*var, active_outputs);
|
|
|
+ for (auto &var : output_variables)
|
|
|
+ {
|
|
|
+ if (var.block)
|
|
|
+ emit_interface_block_member_in_struct(*var.var, var.block_member_index, var.location, active_outputs);
|
|
|
+ else
|
|
|
+ emit_interface_block_in_struct(*var.var, active_outputs);
|
|
|
+ }
|
|
|
emit_builtin_outputs_in_struct();
|
|
|
end_scope_decl();
|
|
|
statement("");
|
|
|
@@ -1944,6 +1940,28 @@ void CompilerHLSL::emit_resources()
|
|
|
end_scope();
|
|
|
statement("");
|
|
|
}
|
|
|
+
|
|
|
+ for (TypeID type_id : composite_selection_workaround_types)
|
|
|
+ {
|
|
|
+ // Need out variable since HLSL does not support returning arrays.
|
|
|
+ auto &type = get<SPIRType>(type_id);
|
|
|
+ auto type_str = type_to_glsl(type);
|
|
|
+ auto type_arr_str = type_to_array_glsl(type);
|
|
|
+ statement("void spvSelectComposite(out ", type_str, " out_value", type_arr_str, ", bool cond, ",
|
|
|
+ type_str, " true_val", type_arr_str, ", ",
|
|
|
+ type_str, " false_val", type_arr_str, ")");
|
|
|
+ begin_scope();
|
|
|
+ statement("if (cond)");
|
|
|
+ begin_scope();
|
|
|
+ statement("out_value = true_val;");
|
|
|
+ end_scope();
|
|
|
+ statement("else");
|
|
|
+ begin_scope();
|
|
|
+ statement("out_value = false_val;");
|
|
|
+ end_scope();
|
|
|
+ end_scope();
|
|
|
+ statement("");
|
|
|
+ }
|
|
|
}
|
|
|
|
|
|
void CompilerHLSL::emit_texture_size_variants(uint64_t variant_mask, const char *vecsize_qualifier, bool uav,
|
|
|
@@ -2051,13 +2069,6 @@ void CompilerHLSL::emit_struct_member(const SPIRType &type, uint32_t member_type
|
|
|
if (index < memb.size())
|
|
|
memberflags = memb[index].decoration_flags;
|
|
|
|
|
|
- string qualifiers;
|
|
|
- bool is_block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) ||
|
|
|
- ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock);
|
|
|
-
|
|
|
- if (is_block)
|
|
|
- qualifiers = to_interpolation_qualifiers(memberflags);
|
|
|
-
|
|
|
string packing_offset;
|
|
|
bool is_push_constant = type.storage == StorageClassPushConstant;
|
|
|
|
|
|
@@ -2072,7 +2083,7 @@ void CompilerHLSL::emit_struct_member(const SPIRType &type, uint32_t member_type
|
|
|
packing_offset = join(" : packoffset(c", offset / 16, packing_swizzle[(offset & 15) >> 2], ")");
|
|
|
}
|
|
|
|
|
|
- statement(layout_for_member(type, index), qualifiers, qualifier,
|
|
|
+ statement(layout_for_member(type, index), qualifier,
|
|
|
variable_decl(membertype, to_member_name(type, index)), packing_offset, ";");
|
|
|
}
|
|
|
|
|
|
@@ -2249,7 +2260,7 @@ void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var)
|
|
|
|
|
|
string CompilerHLSL::to_sampler_expression(uint32_t id)
|
|
|
{
|
|
|
- auto expr = join("_", to_expression(id));
|
|
|
+ auto expr = join("_", to_non_uniform_aware_expression(id));
|
|
|
auto index = expr.find_first_of('[');
|
|
|
if (index == string::npos)
|
|
|
{
|
|
|
@@ -2369,7 +2380,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret
|
|
|
arg_type.image.dim != DimBuffer)
|
|
|
{
|
|
|
// Manufacture automatic sampler arg for SampledImage texture
|
|
|
- arglist.push_back(join(image_is_comparison(arg_type, arg.id) ? "SamplerComparisonState " : "SamplerState ",
|
|
|
+ arglist.push_back(join(is_depth_image(arg_type, arg.id) ? "SamplerComparisonState " : "SamplerState ",
|
|
|
to_sampler_expression(arg.id), type_to_array_glsl(arg_type)));
|
|
|
}
|
|
|
|
|
|
@@ -2407,27 +2418,6 @@ void CompilerHLSL::emit_hlsl_entry_point()
|
|
|
if (require_input)
|
|
|
arguments.push_back("SPIRV_Cross_Input stage_input");
|
|
|
|
|
|
- // Add I/O blocks as separate arguments with appropriate storage qualifier.
|
|
|
- ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
- auto &type = this->get<SPIRType>(var.basetype);
|
|
|
- bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
|
|
|
-
|
|
|
- if (var.storage != StorageClassInput && var.storage != StorageClassOutput)
|
|
|
- return;
|
|
|
-
|
|
|
- if (block && !is_builtin_variable(var) && interface_variable_exists_in_entry_point(var.self))
|
|
|
- {
|
|
|
- if (var.storage == StorageClassInput)
|
|
|
- {
|
|
|
- arguments.push_back(join("in ", variable_decl(type, join("stage_input", to_name(var.self)))));
|
|
|
- }
|
|
|
- else if (var.storage == StorageClassOutput)
|
|
|
- {
|
|
|
- arguments.push_back(join("out ", variable_decl(type, join("stage_output", to_name(var.self)))));
|
|
|
- }
|
|
|
- }
|
|
|
- });
|
|
|
-
|
|
|
auto &execution = get_entry_point();
|
|
|
|
|
|
switch (execution.model)
|
|
|
@@ -2588,36 +2578,43 @@ void CompilerHLSL::emit_hlsl_entry_point()
|
|
|
// Copy from stage input struct to globals.
|
|
|
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
auto &type = this->get<SPIRType>(var.basetype);
|
|
|
- bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
|
|
|
+ bool block = has_decoration(type.self, DecorationBlock);
|
|
|
|
|
|
if (var.storage != StorageClassInput)
|
|
|
return;
|
|
|
|
|
|
bool need_matrix_unroll = var.storage == StorageClassInput && execution.model == ExecutionModelVertex;
|
|
|
|
|
|
- if (!block && !var.remapped_variable && type.pointer && !is_builtin_variable(var) &&
|
|
|
+ if (!var.remapped_variable && type.pointer && !is_builtin_variable(var) &&
|
|
|
interface_variable_exists_in_entry_point(var.self))
|
|
|
{
|
|
|
- auto name = to_name(var.self);
|
|
|
- auto &mtype = this->get<SPIRType>(var.basetype);
|
|
|
- if (need_matrix_unroll && mtype.columns > 1)
|
|
|
+ if (block)
|
|
|
{
|
|
|
- // Unroll matrices.
|
|
|
- for (uint32_t col = 0; col < mtype.columns; col++)
|
|
|
- statement(name, "[", col, "] = stage_input.", name, "_", col, ";");
|
|
|
+ auto type_name = to_name(type.self);
|
|
|
+ auto var_name = to_name(var.self);
|
|
|
+ for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(type.member_types.size()); mbr_idx++)
|
|
|
+ {
|
|
|
+ auto mbr_name = to_member_name(type, mbr_idx);
|
|
|
+ auto flat_name = join(type_name, "_", mbr_name);
|
|
|
+ statement(var_name, ".", mbr_name, " = stage_input.", flat_name, ";");
|
|
|
+ }
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
- statement(name, " = stage_input.", name, ";");
|
|
|
+ auto name = to_name(var.self);
|
|
|
+ auto &mtype = this->get<SPIRType>(var.basetype);
|
|
|
+ if (need_matrix_unroll && mtype.columns > 1)
|
|
|
+ {
|
|
|
+ // Unroll matrices.
|
|
|
+ for (uint32_t col = 0; col < mtype.columns; col++)
|
|
|
+ statement(name, "[", col, "] = stage_input.", name, "_", col, ";");
|
|
|
+ }
|
|
|
+ else
|
|
|
+ {
|
|
|
+ statement(name, " = stage_input.", name, ";");
|
|
|
+ }
|
|
|
}
|
|
|
}
|
|
|
-
|
|
|
- // I/O blocks don't use the common stage input/output struct, but separate outputs.
|
|
|
- if (block && !is_builtin_variable(var) && interface_variable_exists_in_entry_point(var.self))
|
|
|
- {
|
|
|
- auto name = to_name(var.self);
|
|
|
- statement(name, " = stage_input", name, ";");
|
|
|
- }
|
|
|
});
|
|
|
|
|
|
// Run the shader.
|
|
|
@@ -2630,22 +2627,6 @@ void CompilerHLSL::emit_hlsl_entry_point()
|
|
|
else
|
|
|
SPIRV_CROSS_THROW("Unsupported shader stage.");
|
|
|
|
|
|
- // Copy block outputs.
|
|
|
- ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
- auto &type = this->get<SPIRType>(var.basetype);
|
|
|
- bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
|
|
|
-
|
|
|
- if (var.storage != StorageClassOutput)
|
|
|
- return;
|
|
|
-
|
|
|
- // I/O blocks don't use the common stage input/output struct, but separate outputs.
|
|
|
- if (block && !is_builtin_variable(var) && interface_variable_exists_in_entry_point(var.self))
|
|
|
- {
|
|
|
- auto name = to_name(var.self);
|
|
|
- statement("stage_output", name, " = ", name, ";");
|
|
|
- }
|
|
|
- });
|
|
|
-
|
|
|
// Copy stage outputs.
|
|
|
if (require_output)
|
|
|
{
|
|
|
@@ -2682,27 +2663,43 @@ void CompilerHLSL::emit_hlsl_entry_point()
|
|
|
|
|
|
ir.for_each_typed_id<SPIRVariable>([&](uint32_t, SPIRVariable &var) {
|
|
|
auto &type = this->get<SPIRType>(var.basetype);
|
|
|
- bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock);
|
|
|
+ bool block = has_decoration(type.self, DecorationBlock);
|
|
|
|
|
|
if (var.storage != StorageClassOutput)
|
|
|
return;
|
|
|
|
|
|
- if (!block && var.storage != StorageClassFunction && !var.remapped_variable && type.pointer &&
|
|
|
- !is_builtin_variable(var) && interface_variable_exists_in_entry_point(var.self))
|
|
|
+ if (!var.remapped_variable && type.pointer &&
|
|
|
+ !is_builtin_variable(var) &&
|
|
|
+ interface_variable_exists_in_entry_point(var.self))
|
|
|
{
|
|
|
- auto name = to_name(var.self);
|
|
|
-
|
|
|
- if (legacy && execution.model == ExecutionModelFragment)
|
|
|
+ if (block)
|
|
|
{
|
|
|
- string output_filler;
|
|
|
- for (uint32_t size = type.vecsize; size < 4; ++size)
|
|
|
- output_filler += ", 0.0";
|
|
|
-
|
|
|
- statement("stage_output.", name, " = float4(", name, output_filler, ");");
|
|
|
+ // I/O blocks need to flatten output.
|
|
|
+ auto type_name = to_name(type.self);
|
|
|
+ auto var_name = to_name(var.self);
|
|
|
+ for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(type.member_types.size()); mbr_idx++)
|
|
|
+ {
|
|
|
+ auto mbr_name = to_member_name(type, mbr_idx);
|
|
|
+ auto flat_name = join(type_name, "_", mbr_name);
|
|
|
+ statement("stage_output.", flat_name, " = ", var_name, ".", mbr_name, ";");
|
|
|
+ }
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
- statement("stage_output.", name, " = ", name, ";");
|
|
|
+ auto name = to_name(var.self);
|
|
|
+
|
|
|
+ if (legacy && execution.model == ExecutionModelFragment)
|
|
|
+ {
|
|
|
+ string output_filler;
|
|
|
+ for (uint32_t size = type.vecsize; size < 4; ++size)
|
|
|
+ output_filler += ", 0.0";
|
|
|
+
|
|
|
+ statement("stage_output.", name, " = float4(", name, output_filler, ");");
|
|
|
+ }
|
|
|
+ else
|
|
|
+ {
|
|
|
+ statement("stage_output.", name, " = ", name, ";");
|
|
|
+ }
|
|
|
}
|
|
|
}
|
|
|
});
|
|
|
@@ -2754,13 +2751,16 @@ void CompilerHLSL::emit_texture_op(const Instruction &i, bool sparse)
|
|
|
bool proj = false;
|
|
|
const uint32_t *opt = nullptr;
|
|
|
auto *combined_image = maybe_get<SPIRCombinedImageSampler>(img);
|
|
|
- auto img_expr = to_expression(combined_image ? combined_image->image : img);
|
|
|
|
|
|
- inherited_expressions.push_back(coord);
|
|
|
+ if (combined_image && has_decoration(img, DecorationNonUniform))
|
|
|
+ {
|
|
|
+ set_decoration(combined_image->image, DecorationNonUniform);
|
|
|
+ set_decoration(combined_image->sampler, DecorationNonUniform);
|
|
|
+ }
|
|
|
+
|
|
|
+ auto img_expr = to_non_uniform_aware_expression(combined_image ? combined_image->image : img);
|
|
|
|
|
|
- // Make sure non-uniform decoration is back-propagated to where it needs to be.
|
|
|
- if (has_decoration(img, DecorationNonUniformEXT))
|
|
|
- propagate_nonuniform_qualifier(img);
|
|
|
+ inherited_expressions.push_back(coord);
|
|
|
|
|
|
switch (op)
|
|
|
{
|
|
|
@@ -2910,7 +2910,7 @@ void CompilerHLSL::emit_texture_op(const Instruction &i, bool sparse)
|
|
|
{
|
|
|
texop += img_expr;
|
|
|
|
|
|
- if (image_is_comparison(imgtype, img))
|
|
|
+ if (is_depth_image(imgtype, img))
|
|
|
{
|
|
|
if (gather)
|
|
|
{
|
|
|
@@ -3016,7 +3016,7 @@ void CompilerHLSL::emit_texture_op(const Instruction &i, bool sparse)
|
|
|
{
|
|
|
string sampler_expr;
|
|
|
if (combined_image)
|
|
|
- sampler_expr = to_expression(combined_image->sampler);
|
|
|
+ sampler_expr = to_non_uniform_aware_expression(combined_image->sampler);
|
|
|
else
|
|
|
sampler_expr = to_sampler_expression(img);
|
|
|
expr += sampler_expr;
|
|
|
@@ -3386,7 +3386,7 @@ void CompilerHLSL::emit_modern_uniform(const SPIRVariable &var)
|
|
|
if (type.basetype == SPIRType::SampledImage && type.image.dim != DimBuffer)
|
|
|
{
|
|
|
// For combined image samplers, also emit a combined image sampler.
|
|
|
- if (image_is_comparison(type, var.self))
|
|
|
+ if (is_depth_image(type, var.self))
|
|
|
statement("SamplerComparisonState ", to_sampler_expression(var.self), type_to_array_glsl(type),
|
|
|
to_resource_binding_sampler(var), ";");
|
|
|
else
|
|
|
@@ -3812,6 +3812,10 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
|
|
|
SPIRV_CROSS_THROW("Reading types other than 32-bit from ByteAddressBuffer not yet supported, unless SM 6.2 and "
|
|
|
"native 16-bit types are enabled.");
|
|
|
|
|
|
+ string base = chain.base;
|
|
|
+ if (has_decoration(chain.self, DecorationNonUniform))
|
|
|
+ convert_non_uniform_expression(base, chain.self);
|
|
|
+
|
|
|
bool templated_load = hlsl_options.shader_model >= 62;
|
|
|
string load_expr;
|
|
|
|
|
|
@@ -3844,7 +3848,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
|
|
|
if (templated_load)
|
|
|
load_op = "Load";
|
|
|
|
|
|
- load_expr = join(chain.base, ".", load_op, template_expr, "(", chain.dynamic_index, chain.static_index, ")");
|
|
|
+ load_expr = join(base, ".", load_op, template_expr, "(", chain.dynamic_index, chain.static_index, ")");
|
|
|
}
|
|
|
else if (type.columns == 1)
|
|
|
{
|
|
|
@@ -3866,7 +3870,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
|
|
|
|
|
|
for (uint32_t r = 0; r < type.vecsize; r++)
|
|
|
{
|
|
|
- load_expr += join(chain.base, ".Load", template_expr, "(", chain.dynamic_index,
|
|
|
+ load_expr += join(base, ".Load", template_expr, "(", chain.dynamic_index,
|
|
|
chain.static_index + r * chain.matrix_stride, ")");
|
|
|
if (r + 1 < type.vecsize)
|
|
|
load_expr += ", ";
|
|
|
@@ -3915,7 +3919,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
|
|
|
|
|
|
for (uint32_t c = 0; c < type.columns; c++)
|
|
|
{
|
|
|
- load_expr += join(chain.base, ".", load_op, template_expr, "(", chain.dynamic_index,
|
|
|
+ load_expr += join(base, ".", load_op, template_expr, "(", chain.dynamic_index,
|
|
|
chain.static_index + c * chain.matrix_stride, ")");
|
|
|
if (c + 1 < type.columns)
|
|
|
load_expr += ", ";
|
|
|
@@ -3944,7 +3948,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR
|
|
|
{
|
|
|
for (uint32_t r = 0; r < type.vecsize; r++)
|
|
|
{
|
|
|
- load_expr += join(chain.base, ".Load", template_expr, "(", chain.dynamic_index,
|
|
|
+ load_expr += join(base, ".Load", template_expr, "(", chain.dynamic_index,
|
|
|
chain.static_index + c * (type.width / 8) + r * chain.matrix_stride, ")");
|
|
|
|
|
|
if ((r + 1 < type.vecsize) || (c + 1 < type.columns))
|
|
|
@@ -3981,9 +3985,6 @@ void CompilerHLSL::emit_load(const Instruction &instruction)
|
|
|
uint32_t id = ops[1];
|
|
|
uint32_t ptr = ops[2];
|
|
|
|
|
|
- if (has_decoration(ptr, DecorationNonUniformEXT))
|
|
|
- propagate_nonuniform_qualifier(ptr);
|
|
|
-
|
|
|
auto &type = get<SPIRType>(result_type);
|
|
|
bool composite_load = !type.array.empty() || type.basetype == SPIRType::Struct;
|
|
|
|
|
|
@@ -4122,9 +4123,6 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|
|
// Make sure we trigger a read of the constituents in the access chain.
|
|
|
track_expression_read(chain.self);
|
|
|
|
|
|
- if (has_decoration(chain.self, DecorationNonUniformEXT))
|
|
|
- propagate_nonuniform_qualifier(chain.self);
|
|
|
-
|
|
|
SPIRType target_type;
|
|
|
target_type.basetype = SPIRType::UInt;
|
|
|
target_type.vecsize = type.vecsize;
|
|
|
@@ -4148,6 +4146,10 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|
|
|
|
|
bool templated_store = hlsl_options.shader_model >= 62;
|
|
|
|
|
|
+ auto base = chain.base;
|
|
|
+ if (has_decoration(chain.self, DecorationNonUniform))
|
|
|
+ convert_non_uniform_expression(base, chain.self);
|
|
|
+
|
|
|
string template_expr;
|
|
|
if (templated_store)
|
|
|
template_expr = join("<", type_to_glsl(type), ">");
|
|
|
@@ -4183,7 +4185,7 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|
|
}
|
|
|
else
|
|
|
store_op = "Store";
|
|
|
- statement(chain.base, ".", store_op, template_expr, "(", chain.dynamic_index, chain.static_index, ", ",
|
|
|
+ statement(base, ".", store_op, template_expr, "(", chain.dynamic_index, chain.static_index, ", ",
|
|
|
store_expr, ");");
|
|
|
}
|
|
|
else if (type.columns == 1)
|
|
|
@@ -4214,7 +4216,7 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|
|
store_expr = join(bitcast_op, "(", store_expr, ")");
|
|
|
}
|
|
|
|
|
|
- statement(chain.base, ".Store", template_expr, "(", chain.dynamic_index,
|
|
|
+ statement(base, ".Store", template_expr, "(", chain.dynamic_index,
|
|
|
chain.static_index + chain.matrix_stride * r, ", ", store_expr, ");");
|
|
|
}
|
|
|
}
|
|
|
@@ -4258,7 +4260,7 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|
|
store_expr = join(bitcast_op, "(", store_expr, ")");
|
|
|
}
|
|
|
|
|
|
- statement(chain.base, ".", store_op, template_expr, "(", chain.dynamic_index,
|
|
|
+ statement(base, ".", store_op, template_expr, "(", chain.dynamic_index,
|
|
|
chain.static_index + c * chain.matrix_stride, ", ", store_expr, ");");
|
|
|
}
|
|
|
}
|
|
|
@@ -4282,7 +4284,7 @@ void CompilerHLSL::write_access_chain(const SPIRAccessChain &chain, uint32_t val
|
|
|
auto bitcast_op = bitcast_glsl_op(target_type, type);
|
|
|
if (!bitcast_op.empty())
|
|
|
store_expr = join(bitcast_op, "(", store_expr, ")");
|
|
|
- statement(chain.base, ".Store", template_expr, "(", chain.dynamic_index,
|
|
|
+ statement(base, ".Store", template_expr, "(", chain.dynamic_index,
|
|
|
chain.static_index + c * (type.width / 8) + r * chain.matrix_stride, ", ", store_expr, ");");
|
|
|
}
|
|
|
}
|
|
|
@@ -4384,9 +4386,6 @@ void CompilerHLSL::emit_access_chain(const Instruction &instruction)
|
|
|
inherit_expression_dependencies(ops[1], ops[i]);
|
|
|
add_implied_read_expression(e, ops[i]);
|
|
|
}
|
|
|
-
|
|
|
- if (has_decoration(ops[1], DecorationNonUniformEXT))
|
|
|
- propagate_nonuniform_qualifier(ops[1]);
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
@@ -4486,13 +4485,16 @@ void CompilerHLSL::emit_atomic(const uint32_t *ops, uint32_t length, spv::Op op)
|
|
|
|
|
|
if (data_type.storage == StorageClassImage || !chain)
|
|
|
{
|
|
|
- statement(atomic_op, "(", to_expression(ops[0]), ", ", to_expression(ops[3]), ", ", to_expression(tmp_id),
|
|
|
- ");");
|
|
|
+ statement(atomic_op, "(", to_non_uniform_aware_expression(ops[0]), ", ",
|
|
|
+ to_expression(ops[3]), ", ", to_expression(tmp_id), ");");
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
+ string base = chain->base;
|
|
|
+ if (has_decoration(chain->self, DecorationNonUniform))
|
|
|
+ convert_non_uniform_expression(base, chain->self);
|
|
|
// RWByteAddress buffer is always uint in its underlying type.
|
|
|
- statement(chain->base, ".", atomic_op, "(", chain->dynamic_index, chain->static_index, ", ",
|
|
|
+ statement(base, ".", atomic_op, "(", chain->dynamic_index, chain->static_index, ", ",
|
|
|
to_expression(ops[3]), ", ", to_expression(tmp_id), ");");
|
|
|
}
|
|
|
}
|
|
|
@@ -4510,14 +4512,17 @@ void CompilerHLSL::emit_atomic(const uint32_t *ops, uint32_t length, spv::Op op)
|
|
|
SPIRType::BaseType expr_type;
|
|
|
if (data_type.storage == StorageClassImage || !chain)
|
|
|
{
|
|
|
- statement(atomic_op, "(", to_expression(ops[2]), ", ", value_expr, ", ", to_name(id), ");");
|
|
|
+ statement(atomic_op, "(", to_non_uniform_aware_expression(ops[2]), ", ", value_expr, ", ", to_name(id), ");");
|
|
|
expr_type = data_type.basetype;
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
// RWByteAddress buffer is always uint in its underlying type.
|
|
|
+ string base = chain->base;
|
|
|
+ if (has_decoration(chain->self, DecorationNonUniform))
|
|
|
+ convert_non_uniform_expression(base, chain->self);
|
|
|
expr_type = SPIRType::UInt;
|
|
|
- statement(chain->base, ".", atomic_op, "(", chain->dynamic_index, chain->static_index, ", ", value_expr,
|
|
|
+ statement(base, ".", atomic_op, "(", chain->dynamic_index, chain->static_index, ", ", value_expr,
|
|
|
", ", to_name(id), ");");
|
|
|
}
|
|
|
|
|
|
@@ -4581,19 +4586,15 @@ void CompilerHLSL::emit_subgroup_op(const Instruction &i)
|
|
|
|
|
|
case OpGroupNonUniformInverseBallot:
|
|
|
SPIRV_CROSS_THROW("Cannot trivially implement InverseBallot in HLSL.");
|
|
|
- break;
|
|
|
|
|
|
case OpGroupNonUniformBallotBitExtract:
|
|
|
SPIRV_CROSS_THROW("Cannot trivially implement BallotBitExtract in HLSL.");
|
|
|
- break;
|
|
|
|
|
|
case OpGroupNonUniformBallotFindLSB:
|
|
|
SPIRV_CROSS_THROW("Cannot trivially implement BallotFindLSB in HLSL.");
|
|
|
- break;
|
|
|
|
|
|
case OpGroupNonUniformBallotFindMSB:
|
|
|
SPIRV_CROSS_THROW("Cannot trivially implement BallotFindMSB in HLSL.");
|
|
|
- break;
|
|
|
|
|
|
case OpGroupNonUniformBallotBitCount:
|
|
|
{
|
|
|
@@ -4618,13 +4619,35 @@ void CompilerHLSL::emit_subgroup_op(const Instruction &i)
|
|
|
}
|
|
|
|
|
|
case OpGroupNonUniformShuffle:
|
|
|
- SPIRV_CROSS_THROW("Cannot trivially implement Shuffle in HLSL.");
|
|
|
+ emit_binary_func_op(result_type, id, ops[3], ops[4], "WaveReadLaneAt");
|
|
|
+ break;
|
|
|
case OpGroupNonUniformShuffleXor:
|
|
|
- SPIRV_CROSS_THROW("Cannot trivially implement ShuffleXor in HLSL.");
|
|
|
+ {
|
|
|
+ bool forward = should_forward(ops[3]);
|
|
|
+ emit_op(ops[0], ops[1],
|
|
|
+ join("WaveReadLaneAt(", to_unpacked_expression(ops[3]), ", ",
|
|
|
+ "WaveGetLaneIndex() ^ ", to_enclosed_expression(ops[4]), ")"), forward);
|
|
|
+ inherit_expression_dependencies(ops[1], ops[3]);
|
|
|
+ break;
|
|
|
+ }
|
|
|
case OpGroupNonUniformShuffleUp:
|
|
|
- SPIRV_CROSS_THROW("Cannot trivially implement ShuffleUp in HLSL.");
|
|
|
+ {
|
|
|
+ bool forward = should_forward(ops[3]);
|
|
|
+ emit_op(ops[0], ops[1],
|
|
|
+ join("WaveReadLaneAt(", to_unpacked_expression(ops[3]), ", ",
|
|
|
+ "WaveGetLaneIndex() - ", to_enclosed_expression(ops[4]), ")"), forward);
|
|
|
+ inherit_expression_dependencies(ops[1], ops[3]);
|
|
|
+ break;
|
|
|
+ }
|
|
|
case OpGroupNonUniformShuffleDown:
|
|
|
- SPIRV_CROSS_THROW("Cannot trivially implement ShuffleDown in HLSL.");
|
|
|
+ {
|
|
|
+ bool forward = should_forward(ops[3]);
|
|
|
+ emit_op(ops[0], ops[1],
|
|
|
+ join("WaveReadLaneAt(", to_unpacked_expression(ops[3]), ", ",
|
|
|
+ "WaveGetLaneIndex() + ", to_enclosed_expression(ops[4]), ")"), forward);
|
|
|
+ inherit_expression_dependencies(ops[1], ops[3]);
|
|
|
+ break;
|
|
|
+ }
|
|
|
|
|
|
case OpGroupNonUniformAll:
|
|
|
emit_unary_func_op(result_type, id, ops[3], "WaveActiveAllTrue");
|
|
|
@@ -4771,6 +4794,34 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
break;
|
|
|
}
|
|
|
|
|
|
+ case OpSelect:
|
|
|
+ {
|
|
|
+ auto &value_type = expression_type(ops[3]);
|
|
|
+ if (value_type.basetype == SPIRType::Struct || is_array(value_type))
|
|
|
+ {
|
|
|
+ // HLSL does not support ternary expressions on composites.
|
|
|
+ // Cannot use branches, since we might be in a continue block
|
|
|
+ // where explicit control flow is prohibited.
|
|
|
+ // Emit a helper function where we can use control flow.
|
|
|
+ TypeID value_type_id = expression_type_id(ops[3]);
|
|
|
+ auto itr = std::find(composite_selection_workaround_types.begin(),
|
|
|
+ composite_selection_workaround_types.end(),
|
|
|
+ value_type_id);
|
|
|
+ if (itr == composite_selection_workaround_types.end())
|
|
|
+ {
|
|
|
+ composite_selection_workaround_types.push_back(value_type_id);
|
|
|
+ force_recompile();
|
|
|
+ }
|
|
|
+ emit_uninitialized_temporary_expression(ops[0], ops[1]);
|
|
|
+ statement("spvSelectComposite(",
|
|
|
+ to_expression(ops[1]), ", ", to_expression(ops[2]), ", ",
|
|
|
+ to_expression(ops[3]), ", ", to_expression(ops[4]), ");");
|
|
|
+ }
|
|
|
+ else
|
|
|
+ CompilerGLSL::emit_instruction(instruction);
|
|
|
+ break;
|
|
|
+ }
|
|
|
+
|
|
|
case OpStore:
|
|
|
{
|
|
|
emit_store(instruction);
|
|
|
@@ -5150,7 +5201,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
auto dummy_samples_levels = join(get_fallback_name(id), "_dummy_parameter");
|
|
|
statement("uint ", dummy_samples_levels, ";");
|
|
|
|
|
|
- auto expr = join("spvTextureSize(", to_expression(ops[2]), ", ",
|
|
|
+ auto expr = join("spvTextureSize(", to_non_uniform_aware_expression(ops[2]), ", ",
|
|
|
bitcast_expression(SPIRType::UInt, ops[3]), ", ", dummy_samples_levels, ")");
|
|
|
|
|
|
auto &restype = get<SPIRType>(ops[0]);
|
|
|
@@ -5176,9 +5227,9 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
|
|
|
string expr;
|
|
|
if (uav)
|
|
|
- expr = join("spvImageSize(", to_expression(ops[2]), ", ", dummy_samples_levels, ")");
|
|
|
+ expr = join("spvImageSize(", to_non_uniform_aware_expression(ops[2]), ", ", dummy_samples_levels, ")");
|
|
|
else
|
|
|
- expr = join("spvTextureSize(", to_expression(ops[2]), ", 0u, ", dummy_samples_levels, ")");
|
|
|
+ expr = join("spvTextureSize(", to_non_uniform_aware_expression(ops[2]), ", 0u, ", dummy_samples_levels, ")");
|
|
|
|
|
|
auto &restype = get<SPIRType>(ops[0]);
|
|
|
expr = bitcast_expression(restype, SPIRType::UInt, expr);
|
|
|
@@ -5208,9 +5259,9 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
statement(variable_decl(type, to_name(id)), ";");
|
|
|
|
|
|
if (uav)
|
|
|
- statement("spvImageSize(", to_expression(ops[2]), ", ", to_name(id), ");");
|
|
|
+ statement("spvImageSize(", to_non_uniform_aware_expression(ops[2]), ", ", to_name(id), ");");
|
|
|
else
|
|
|
- statement("spvTextureSize(", to_expression(ops[2]), ", 0u, ", to_name(id), ");");
|
|
|
+ statement("spvTextureSize(", to_non_uniform_aware_expression(ops[2]), ", 0u, ", to_name(id), ");");
|
|
|
|
|
|
auto &restype = get<SPIRType>(ops[0]);
|
|
|
auto expr = bitcast_expression(restype, SPIRType::UInt, to_name(id));
|
|
|
@@ -5241,16 +5292,16 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
if (operands != ImageOperandsSampleMask || instruction.length != 6)
|
|
|
SPIRV_CROSS_THROW("Multisampled image used in OpImageRead, but unexpected operand mask was used.");
|
|
|
uint32_t sample = ops[5];
|
|
|
- imgexpr = join(to_expression(ops[2]), ".Load(int2(gl_FragCoord.xy), ", to_expression(sample), ")");
|
|
|
+ imgexpr = join(to_non_uniform_aware_expression(ops[2]), ".Load(int2(gl_FragCoord.xy), ", to_expression(sample), ")");
|
|
|
}
|
|
|
else
|
|
|
- imgexpr = join(to_expression(ops[2]), ".Load(int3(int2(gl_FragCoord.xy), 0))");
|
|
|
+ imgexpr = join(to_non_uniform_aware_expression(ops[2]), ".Load(int3(int2(gl_FragCoord.xy), 0))");
|
|
|
|
|
|
pure = true;
|
|
|
}
|
|
|
else
|
|
|
{
|
|
|
- imgexpr = join(to_expression(ops[2]), "[", to_expression(ops[3]), "]");
|
|
|
+ imgexpr = join(to_non_uniform_aware_expression(ops[2]), "[", to_expression(ops[3]), "]");
|
|
|
// The underlying image type in HLSL depends on the image format, unlike GLSL, where all images are "vec4",
|
|
|
// except that the underlying type changes how the data is interpreted.
|
|
|
|
|
|
@@ -5299,7 +5350,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
value_expr = remap_swizzle(narrowed_type, expression_type(ops[2]).vecsize, value_expr);
|
|
|
}
|
|
|
|
|
|
- statement(to_expression(ops[0]), "[", to_expression(ops[1]), "] = ", value_expr, ";");
|
|
|
+ statement(to_non_uniform_aware_expression(ops[0]), "[", to_expression(ops[1]), "] = ", value_expr, ";");
|
|
|
if (var && variable_storage_is_aliased(*var))
|
|
|
flush_all_aliased_variables();
|
|
|
break;
|
|
|
@@ -5311,10 +5362,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
uint32_t id = ops[1];
|
|
|
|
|
|
auto expr = to_expression(ops[2]);
|
|
|
- if (has_decoration(id, DecorationNonUniformEXT) || has_decoration(ops[2], DecorationNonUniformEXT))
|
|
|
- convert_non_uniform_expression(expression_type(ops[2]), expr);
|
|
|
expr += join("[", to_expression(ops[3]), "]");
|
|
|
-
|
|
|
auto &e = set<SPIRExpression>(id, expr, result_type, true);
|
|
|
|
|
|
// When using the pointer, we need to know which variable it is actually loaded from.
|
|
|
@@ -5492,7 +5540,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
|
|
|
case OpArrayLength:
|
|
|
{
|
|
|
- auto *var = maybe_get<SPIRVariable>(ops[2]);
|
|
|
+ auto *var = maybe_get_backing_variable(ops[2]);
|
|
|
if (!var)
|
|
|
SPIRV_CROSS_THROW("Array length must point directly to an SSBO block.");
|
|
|
|
|
|
@@ -5502,7 +5550,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction)
|
|
|
|
|
|
// This must be 32-bit uint, so we're good to go.
|
|
|
emit_uninitialized_temporary_expression(ops[0], ops[1]);
|
|
|
- statement(to_expression(ops[2]), ".GetDimensions(", to_expression(ops[1]), ");");
|
|
|
+ statement(to_non_uniform_aware_expression(ops[2]), ".GetDimensions(", to_expression(ops[1]), ");");
|
|
|
uint32_t offset = type_struct_member_offset(type, ops[3]);
|
|
|
uint32_t stride = type_struct_member_array_stride(type, ops[3]);
|
|
|
statement(to_expression(ops[1]), " = (", to_expression(ops[1]), " - ", offset, ") / ", stride, ";");
|
|
|
@@ -5648,6 +5696,7 @@ VariableID CompilerHLSL::remap_num_workgroups_builtin()
|
|
|
ir.meta[variable_id].decoration.alias = "SPIRV_Cross_NumWorkgroups";
|
|
|
|
|
|
num_workgroups_builtin = variable_id;
|
|
|
+ get_entry_point().interface_variables.push_back(num_workgroups_builtin);
|
|
|
return variable_id;
|
|
|
}
|
|
|
|
|
|
@@ -5718,6 +5767,9 @@ string CompilerHLSL::compile()
|
|
|
backend.nonuniform_qualifier = "NonUniformResourceIndex";
|
|
|
backend.support_case_fallthrough = false;
|
|
|
|
|
|
+ // SM 4.1 does not support precise for some reason.
|
|
|
+ backend.support_precise_qualifier = hlsl_options.shader_model >= 50 || hlsl_options.shader_model == 40;
|
|
|
+
|
|
|
fixup_type_alias();
|
|
|
reorder_type_alias();
|
|
|
build_function_control_flow_graphs_and_analyze();
|
Alex Szpakowski