Working code

Author: christiangnrd

src/device/intrinsics/math.jl

@@ -298,14 +298,14 @@ end
     if metal_version() >= sv"3.1" # macOS 14+
         ccall("extern air.nextafter.f32", llvmcall, Cfloat, (Cfloat, Cfloat), x, y)
     else
-        reinterpret(Float32, reinterpret(UInt32, x) + sign(y-x))
+        nextfloat(x, unsafe_trunc(Int32, sign(y - x)))
     end
 end
 @device_function function nextafter(x::Float16, y::Float16)
     if metal_version() >= sv"3.1" # macOS 14+
         ccall("extern air.nextafter.f16", llvmcall, Float16, (Float16, Float16), x, y)
     else
-        reinterpret(Float16, reinterpret(UInt16, x) + sign(y-x))
+        nextfloat(x, unsafe_trunc(Int16, sign(y - x)))
     end
 end
 

test/device/intrinsics.jl

@@ -355,21 +355,19 @@ end
 
 
     let # nextafter
-        if Metal.is_macos(v"14")
-            N = 4
-            function nextafter_test(X, y)
-                idx = thread_position_in_grid_1d()
-                X[idx] = Metal.nextafter(X[idx], y)
-                return nothing
-            end
-            arr = rand(T, N)
-            buffer = MtlArray(arr)
-            Metal.@sync @metal threads = N nextafter_test(buffer, typemax(T))
-            @test Array(buffer) == nextfloat.(arr)
-
-            Metal.@sync @metal threads = N nextafter_test(buffer, typemin(T))
-            @test Array(buffer) == arr
+        N = 4
+        function nextafter_test(X, y)
+            idx = thread_position_in_grid_1d()
+            X[idx] = Metal.nextafter(X[idx], y)
+            return nothing
         end
+        arr = rand(T, N)
+        buffer = MtlArray(arr)
+        Metal.@sync @metal threads = N nextafter_test(buffer, typemax(T))
+        @test Array(buffer) == nextfloat.(arr)
+
+        Metal.@sync @metal threads = N nextafter_test(buffer, typemin(T))
+        @test Array(buffer) == arr
     end
 end
 end