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

MSL: Also replace bool with short in structures.

Browse Source 
Since `bool` is a logical type, it cannot be used in uniform or storage
buffers. Therefore, replacing it in structures should not change the
shader interface.

We leave it alone for builtins. (FIXME: Should we also leave it for
I/O varyings?)

Fixes 24 CTS tests under `dEQP-VK.memory_model.shared`.
This commit is contained in:
Chip Davis 2022-07-23 23:12:26 -07:00
parent 0cccd0a65a
commit faea931de3
5 changed files with 230 additions and 8 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

63
reference/opt/shaders-msl/comp/shared-struct-bool-cast.comp Normal file
View File

@ -0,0 +1,63 @@
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct S1
{
int3 a;
uint2 b;
short4 c;
uint d;
};
struct block
{
uint passed;
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
kernel void main0(device block& _132 [[buffer(0)]])
{
threadgroup S1 s1;
s1.a = int3(6, 8, 8);
s1.b = uint2(4u);
s1.c = short4(bool4(false, false, false, true));
s1.d = 6u;
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
bool _144 = all(int3(6, 8, 8) == s1.a);
bool _108;
if (_144)
{
_108 = all(uint2(4u) == s1.b);
}
else
{
_108 = _144;
}
bool _117;
if (_108)
{
_117 = all(bool4(false, false, false, true) == bool4(s1.c));
}
else
{
_117 = _108;
}
bool _126;
if (_117)
{
_126 = 6u == s1.d;
}
else
{
_126 = _117;
}
if (_126)
{
_132.passed++;
}
}

110
reference/shaders-msl/comp/shared-struct-bool-cast.comp Normal file
View File

@ -0,0 +1,110 @@
#pragma clang diagnostic ignored "-Wmissing-prototypes"
#include <metal_stdlib>
#include <simd/simd.h>
using namespace metal;
struct S1
{
int3 a;
uint2 b;
short4 c;
uint d;
};
struct block
{
uint passed;
};
constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
static inline __attribute__((always_inline))
bool compare_ivec3(thread const int3& a, thread const int3& b)
{
return all(a == b);
}
static inline __attribute__((always_inline))
bool compare_uvec2(thread const uint2& a, thread const uint2& b)
{
return all(a == b);
}
static inline __attribute__((always_inline))
bool compare_bvec4(thread const bool4& a, thread const bool4& b)
{
return all(a == b);
}
static inline __attribute__((always_inline))
bool compare_uint(thread const uint& a, thread const uint& b)
{
return a == b;
}
kernel void main0(device block& _132 [[buffer(0)]])
{
threadgroup S1 s1;
s1.a = int3(6, 8, 8);
s1.b = uint2(4u);
s1.c = short4(bool4(false, false, false, true));
s1.d = 6u;
threadgroup_barrier(mem_flags::mem_threadgroup);
threadgroup_barrier(mem_flags::mem_device | mem_flags::mem_threadgroup | mem_flags::mem_texture);
bool allOk = true;
bool _99;
if (allOk)
{
int3 param = int3(6, 8, 8);
int3 param_1 = s1.a;
_99 = compare_ivec3(param, param_1);
}
else
{
_99 = allOk;
}
allOk = _99;
bool _108;
if (allOk)
{
uint2 param_2 = uint2(4u);
uint2 param_3 = s1.b;
_108 = compare_uvec2(param_2, param_3);
}
else
{
_108 = allOk;
}
allOk = _108;
bool _117;
if (allOk)
{
bool4 param_4 = bool4(false, false, false, true);
bool4 param_5 = bool4(s1.c);
_117 = compare_bvec4(param_4, param_5);
}
else
{
_117 = allOk;
}
allOk = _117;
bool _126;
if (allOk)
{
uint param_6 = 6u;
uint param_7 = s1.d;
_126 = compare_uint(param_6, param_7);
}
else
{
_126 = allOk;
}
allOk = _126;
if (allOk)
{
_132.passed++;
}
}

35
shaders-msl/comp/shared-struct-bool-cast.comp Normal file
View File

@ -0,0 +1,35 @@
#version 450
layout(local_size_x = 1) in;
layout(std140, binding = 0) buffer block { highp uint passed; };
struct S1 {
mediump ivec3 a;
highp uvec2 b;
bvec4 c;
mediump uint d;
};
bool compare_ivec3 (highp ivec3 a, highp ivec3 b) { return a == b; }
bool compare_uint (highp uint a, highp uint b) { return a == b; }
bool compare_uvec2 (highp uvec2 a, highp uvec2 b) { return a == b; }
bool compare_bvec4 (bvec4 a, bvec4 b) { return a == b; }
shared S1 s1;
void main (void) {
s1.a = ivec3(6, 8, 8);
s1.b = uvec2(4u, 4u);
s1.c = bvec4(false, false, false, true);
s1.d = 6u;
barrier();
memoryBarrier();
bool allOk = true;
allOk = allOk && compare_ivec3(ivec3(6, 8, 8), s1.a);
allOk = allOk && compare_uvec2(uvec2(4u, 4u), s1.b);
allOk = allOk && compare_bvec4(bvec4(false, false, false, true), s1.c);
allOk = allOk && compare_uint(6u, s1.d);
if (allOk)
passed++;
}

29
spirv_msl.cpp
View File

@ -10907,8 +10907,8 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_
array_type = type_to_array_glsl(physical_type);
}
auto result = join(pack_pfx, type_to_glsl(*declared_type, orig_id), " ", qualifier, to_member_name(type, index),
member_attribute_qualifier(type, index), array_type, ";");
auto result = join(pack_pfx, type_to_glsl(*declared_type, orig_id, true), " ", qualifier,
to_member_name(type, index), member_attribute_qualifier(type, index), array_type, ";");
is_using_builtin_array = false;
return result;
@ -13581,7 +13581,7 @@ string CompilerMSL::to_qualifiers_glsl(uint32_t id)
// The optional id parameter indicates the object whose type we are trying
// to find the description for. It is optional. Most type descriptions do not
// depend on a specific object's use of that type.
string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member)
{
string type_name;
@ -13671,9 +13671,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
// Need to special-case threadgroup booleans. They are supposed to be logical
// storage, but MSL compilers will sometimes crash if you use threadgroup bool.
// Workaround this by using 16-bit types instead and fixup on load-store to this data.
// FIXME: We have no sane way of working around this problem if a struct member is boolean
// and that struct is used as a threadgroup variable, but ... sigh.
if ((var && var->storage == StorageClassWorkgroup) || type.storage == StorageClassWorkgroup)
if ((var && var->storage == StorageClassWorkgroup) || type.storage == StorageClassWorkgroup || member)
type_name = "short";
else
type_name = "bool";
@ -13765,6 +13763,11 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id)
}
}
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)
{
// Allow Metal to use the array<T> template to make arrays a value type
@ -15787,11 +15790,16 @@ void CompilerMSL::remap_constexpr_sampler_by_binding(uint32_t desc_set, uint32_t
void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type)
{
auto *var = maybe_get_backing_variable(source_id);
SPIRType *var_type;
if (var)
{
source_id = var->self;
var_type = &get_variable_data_type(*var);
}
// Type fixups for workgroup variables if they are booleans.
if (var && var->storage == StorageClassWorkgroup && expr_type.basetype == SPIRType::Boolean)
if (var && (var->storage == StorageClassWorkgroup || var_type->basetype == SPIRType::Struct) &&
expr_type.basetype == SPIRType::Boolean)
expr = join(type_to_glsl(expr_type), "(", expr, ")");
// Only interested in standalone builtin variables in the switch below.
@ -15886,11 +15894,16 @@ void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr,
void CompilerMSL::cast_to_variable_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type)
{
auto *var = maybe_get_backing_variable(target_id);
SPIRType *var_type;
if (var)
{
target_id = var->self;
var_type = &get_variable_data_type(*var);
}
// Type fixups for workgroup variables if they are booleans.
if (var && var->storage == StorageClassWorkgroup && expr_type.basetype == SPIRType::Boolean)
if (var && (var->storage == StorageClassWorkgroup || var_type->basetype == SPIRType::Struct) &&
expr_type.basetype == SPIRType::Boolean)
{
auto short_type = expr_type;
short_type.basetype = SPIRType::Short;

1
spirv_msl.hpp
View File

@ -736,6 +736,7 @@ protected:
void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index,
const std::string &qualifier = "", uint32_t base_offset = 0) override;
void emit_struct_padding_target(const SPIRType &type) override;
std::string type_to_glsl(const SPIRType &type, uint32_t id, bool member);
std::string type_to_glsl(const SPIRType &type, uint32_t id = 0) override;
void emit_block_hints(const SPIRBlock &block) override;