CompilerMSL support platform semantics.
Support customizing MSL based on iOS or macOS platform. Support SPIV-V containing multiple memory semantics.
This commit is contained in:
parent
bcc96d8c70
commit
3fcdce08ab
@ -1559,34 +1559,18 @@ void CompilerMSL::emit_barrier(uint32_t id_exe_scope, uint32_t id_mem_scope, uin
|
|||||||
string bar_stmt = "threadgroup_barrier(mem_flags::";
|
string bar_stmt = "threadgroup_barrier(mem_flags::";
|
||||||
|
|
||||||
uint32_t mem_sem = id_mem_sem ? get<SPIRConstant>(id_mem_sem).scalar() : uint32_t(MemorySemanticsMaskNone);
|
uint32_t mem_sem = id_mem_sem ? get<SPIRConstant>(id_mem_sem).scalar() : uint32_t(MemorySemanticsMaskNone);
|
||||||
switch (mem_sem)
|
|
||||||
{
|
if (mem_sem & MemorySemanticsCrossWorkgroupMemoryMask)
|
||||||
case MemorySemanticsCrossWorkgroupMemoryMask:
|
|
||||||
bar_stmt += "mem_device";
|
bar_stmt += "mem_device";
|
||||||
break;
|
else if (mem_sem & (MemorySemanticsSubgroupMemoryMask | MemorySemanticsWorkgroupMemoryMask |
|
||||||
|
MemorySemanticsAtomicCounterMemoryMask))
|
||||||
case MemorySemanticsSubgroupMemoryMask:
|
|
||||||
case MemorySemanticsWorkgroupMemoryMask:
|
|
||||||
case MemorySemanticsAtomicCounterMemoryMask:
|
|
||||||
bar_stmt += "mem_threadgroup";
|
bar_stmt += "mem_threadgroup";
|
||||||
break;
|
else if (mem_sem & MemorySemanticsImageMemoryMask)
|
||||||
|
|
||||||
case MemorySemanticsImageMemoryMask:
|
|
||||||
bar_stmt += "mem_texture";
|
bar_stmt += "mem_texture";
|
||||||
break;
|
else
|
||||||
|
|
||||||
case MemorySemanticsAcquireMask:
|
|
||||||
case MemorySemanticsReleaseMask:
|
|
||||||
case MemorySemanticsAcquireReleaseMask:
|
|
||||||
case MemorySemanticsSequentiallyConsistentMask:
|
|
||||||
case MemorySemanticsUniformMemoryMask:
|
|
||||||
case MemorySemanticsMaskNone:
|
|
||||||
default:
|
|
||||||
bar_stmt += "mem_none";
|
bar_stmt += "mem_none";
|
||||||
break;
|
|
||||||
}
|
|
||||||
|
|
||||||
if (options.supports_msl_version(2))
|
if (options.is_ios() && options.supports_msl_version(2))
|
||||||
{
|
{
|
||||||
bar_stmt += ", ";
|
bar_stmt += ", ";
|
||||||
|
|
||||||
|
@ -77,10 +77,25 @@ public:
|
|||||||
// Options for compiling to Metal Shading Language
|
// Options for compiling to Metal Shading Language
|
||||||
struct Options
|
struct Options
|
||||||
{
|
{
|
||||||
|
typedef enum {
|
||||||
|
iOS,
|
||||||
|
macOS,
|
||||||
|
} Platform;
|
||||||
|
|
||||||
|
Platform platform = macOS;
|
||||||
uint32_t msl_version = make_msl_version(1, 2);
|
uint32_t msl_version = make_msl_version(1, 2);
|
||||||
bool enable_point_size_builtin = true;
|
bool enable_point_size_builtin = true;
|
||||||
bool resolve_specialized_array_lengths = true;
|
bool resolve_specialized_array_lengths = true;
|
||||||
|
|
||||||
|
bool is_ios()
|
||||||
|
{
|
||||||
|
return platform == iOS;
|
||||||
|
}
|
||||||
|
bool is_macos()
|
||||||
|
{
|
||||||
|
return platform == macOS;
|
||||||
|
}
|
||||||
|
|
||||||
void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0)
|
void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0)
|
||||||
{
|
{
|
||||||
msl_version = make_msl_version(major, minor, patch);
|
msl_version = make_msl_version(major, minor, patch);
|
||||||
|
Loading…
Reference in New Issue
Block a user