Merge pull request #2105 from goki/main

add support for OpAtomicFAddEXT atomic add, and also Min, Max

reference/opt/shaders-msl/comp/atomic-float.msl3.comp

@@ -0,0 +1,20 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float f32;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device SSBO& ssbo [[buffer(0)]])
+{
+    float _18 = atomic_fetch_add_explicit((device atomic_float*)&ssbo.f32, 1.0, memory_order_relaxed);
+}
+

reference/shaders-msl/comp/atomic-float.msl3.comp

@@ -0,0 +1,20 @@
+#pragma clang diagnostic ignored "-Wunused-variable"
+
+#include <metal_stdlib>
+#include <simd/simd.h>
+#include <metal_atomic>
+
+using namespace metal;
+
+struct SSBO
+{
+    float f32;
+};
+
+constant uint3 gl_WorkGroupSize [[maybe_unused]] = uint3(1u);
+
+kernel void main0(device SSBO& ssbo [[buffer(0)]])
+{
+    float _18 = atomic_fetch_add_explicit((device atomic_float*)&ssbo.f32, 1.0, memory_order_relaxed);
+}
+

shaders-msl/comp/atomic-float.msl3.comp

@@ -0,0 +1,15 @@
+#version 450
+#extension GL_EXT_shader_atomic_float : require
+layout(local_size_x = 1) in;
+
+layout(binding = 2, std430) buffer SSBO
+{
+    float f32;
+} ssbo;
+
+void main()
+{
+	// note: only works on ssbo targets
+    atomicAdd(ssbo.f32, 1.0);
+}
+

spirv_msl.cpp

@@ -1856,6 +1856,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std::
 			case OpAtomicIIncrement:
 			case OpAtomicIDecrement:
 			case OpAtomicIAdd:
+			case OpAtomicFAddEXT:
 			case OpAtomicISub:
 			case OpAtomicSMin:
 			case OpAtomicUMin:
@@ -8607,6 +8608,7 @@ void CompilerMSL::emit_instruction(const Instruction &instruction)
 		break;
 
 	case OpAtomicIAdd:
+	case OpAtomicFAddEXT:
 		MSL_AFMO(add);
 		break;
 
@@ -16314,6 +16316,7 @@ bool CompilerMSL::OpCodePreprocessor::handle(Op opcode, const uint32_t *args, ui
 	case OpAtomicIIncrement:
 	case OpAtomicIDecrement:
 	case OpAtomicIAdd:
+	case OpAtomicFAddEXT:
 	case OpAtomicISub:
 	case OpAtomicSMin:
 	case OpAtomicUMin:
@@ -16519,6 +16522,7 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o
 	case OpAtomicIIncrement:
 	case OpAtomicIDecrement:
 	case OpAtomicIAdd:
+	case OpAtomicFAddEXT:
 	case OpAtomicISub:
 	case OpAtomicSMin:
 	case OpAtomicUMin:

test_shaders.py

@@ -111,8 +111,8 @@ def print_msl_compiler_version():
 
 def msl_compiler_supports_version(version):
     try:
-        subprocess.check_call(['xcrun', '--sdk', 'macosx', 'metal', '-x', 'metal', '-std=macos-metal' + version, '-'],
+        subprocess.check_call(['xcrun', '--sdk', 'macosx', 'metal', '-x', 'metal', version, '-'],
-                stdin = subprocess.DEVNULL, stdout = subprocess.DEVNULL, stderr = subprocess.DEVNULL)
+            stdin = subprocess.DEVNULL, stdout = subprocess.DEVNULL, stderr = subprocess.DEVNULL)
         print('Current SDK supports MSL {0}. Enabling validation for MSL {0} shaders.'.format(version))
         return True
     except OSError as e:
@@ -124,7 +124,9 @@ def msl_compiler_supports_version(version):
 
 def path_to_msl_standard(shader):
     if '.ios.' in shader:
-        if '.msl2.' in shader:
+        if '.msl3.' in shader:
+            return '-std=metal3.0'
+        elif '.msl2.' in shader:
             return '-std=ios-metal2.0'
         elif '.msl21.' in shader:
             return '-std=ios-metal2.1'
@@ -141,7 +143,9 @@ def path_to_msl_standard(shader):
         else:
             return '-std=ios-metal1.2'
     else:
-        if '.msl2.' in shader:
+        if '.msl3.' in shader:
+            return '-std=metal3.0'
+        elif '.msl2.' in shader:
             return '-std=macos-metal2.0'
         elif '.msl21.' in shader:
             return '-std=macos-metal2.1'
@@ -157,7 +161,9 @@ def path_to_msl_standard(shader):
             return '-std=macos-metal1.2'
 
 def path_to_msl_standard_cli(shader):
-    if '.msl2.' in shader:
+    if '.msl3.' in shader:
+        return '30000'
+    elif '.msl2.' in shader:
         return '20000'
     elif '.msl21.' in shader:
         return '20100'
@@ -1017,11 +1023,13 @@ def main():
     args.msl22 = False
     args.msl23 = False
     args.msl24 = False
+    args.msl30 = False
     if args.msl:
         print_msl_compiler_version()
-        args.msl22 = msl_compiler_supports_version('2.2')
+        args.msl22 = msl_compiler_supports_version('-std=macos-metal2.2')
-        args.msl23 = msl_compiler_supports_version('2.3')
+        args.msl23 = msl_compiler_supports_version('-std=macos-metal2.3')
-        args.msl24 = msl_compiler_supports_version('2.4')
+        args.msl24 = msl_compiler_supports_version('-std=macos-metal2.4')
+        args.msl30 = msl_compiler_supports_version('-std=metal3.0')
 
     backend = 'glsl'
     if (args.msl or args.metal):