@@ -355,20 +355,33 @@ end
355355
356356
357357 let # nextafter
358- N = 4
359358 function nextafter_test (X, y)
360359 idx = thread_position_in_grid_1d ()
361360 X[idx] = Metal. nextafter (X[idx], y)
362361 return nothing
363362 end
363+
364+ # Check the code is generated as expected
365+ outval = T (0 )
366+ function nextafter_out_test ()
367+ Metal. nextafter (outval, outval)
368+ return
369+ end
370+
371+ N = 4
364372 arr = rand (T, N)
365373
366- # test the intrinsic
367- buffer1 = MtlArray (arr)
368- Metal. @sync @metal threads = N nextafter_test (buffer1, typemax (T))
369- @test Array (buffer1) == nextfloat .(arr)
370- Metal. @sync @metal threads = N nextafter_test (buffer1, typemin (T))
371- @test Array (buffer1) == arr
374+ # test the intrinsic (macOS >= v14)
375+ if metal_support () >= v " 3.1"
376+ buffer1 = MtlArray (arr)
377+ Metal. @sync @metal threads = N nextafter_test (buffer1, typemax (T))
378+ @test Array (buffer1) == nextfloat .(arr)
379+ Metal. @sync @metal threads = N nextafter_test (buffer1, typemin (T))
380+ @test Array (buffer1) == arr
381+
382+ ir = sprint (io-> (@device_code_llvm io= io dump_module= true @metal nextafter_out_test ()))
383+ @test occursin (Regex (" @air\\ .nextafter\\ .f$(8 * sizeof (T)) " ), ir)
384+ end
372385
373386 # test for metal < 3.1
374387 buffer2 = MtlArray (arr)
377390 Metal. @sync @metal threads = N metal = v " 3.0" nextafter_test (buffer2, typemin (T))
378391 @test Array (buffer2) == arr
379392
380- # Check the code is generated as expected
381- outval = T (0 )
382- function nextafter_out_test ()
383- Metal. nextafter (outval, outval)
384- return
385- end
386- ir = sprint (io-> (@device_code_llvm io= io dump_module= true @metal nextafter_out_test ()))
387- @test occursin (Regex (" @air\\ .nextafter\\ .f$(8 * sizeof (T)) " ), ir)
388393 ir = sprint (io-> (@device_code_llvm io= io dump_module= true @metal metal = v " 3.0" nextafter_out_test ()))
389394 @test occursin (Regex (" @air\\ .sign\\ .f$(8 * sizeof (T)) " ), ir)
390395 end
0 commit comments