Test both simd shuffle intrinsics. (#553)

* Docstring fixup

* Test both up and down shuffles

Author: christiangnrd

src/device/intrinsics/simd.jl

@@ -112,10 +112,10 @@ end
 @doc """
     simd_shuffle_down(data::T, delta::Integer)
 
-Return data from the thread whose SIMD lane ID is the sum of caller’s SIMD lane ID and delta.
+Return `data` from the thread whose SIMD lane ID is the sum of caller's SIMD lane ID and `delta`.
 
-The value for delta must be the same for all threads in the SIMD-group. This function
-doesn’t modify the upper delta lanes of data because it doesn’t wrap values around
+The value for `delta` must be the same for all threads in the SIMD-group. This function
+doesn't modify the upper `delta` lanes of `data` because it doesn't wrap values around
 the SIMD-group.
 
 T must be one of the following: Float32, Float16, Int32, UInt32, Int16, UInt16, Int8, or UInt8
@@ -125,11 +125,11 @@ simd_shuffle_down
 @doc """
     simd_shuffle_up(data::T, delta::Integer)
 
-Return data from the thread whose SIMD lane ID is the difference from the caller’s SIMD
-lane ID minus delta.
+Return `data` from the thread whose SIMD lane ID is the difference from the caller's SIMD
+lane ID minus `delta`.
 
-The value of delta must be the same for all threads in a SIMD-group. This function doesn’t
-modify the lower delta lanes of data because it doesn’t wrap values around the SIMD-group.
+The value of `delta` must be the same for all threads in a SIMD-group. This function doesn't
+modify the lower `delta` lanes of `data` because it doesn't wrap values around the SIMD-group.
 
 T must be one of the following: Float32, Float16, Int32, UInt32, Int16, UInt16, Int8, or UInt8
 """

test/device/intrinsics.jl

@@ -625,7 +625,7 @@ end
 
 @testset "simd intrinsics" begin
 
-@testset "shuffle($typ)" for typ in [Float32, Float16, Int32, UInt32, Int16, UInt16, Int8, UInt8]
+@testset "$f($typ)" for typ in [Float32, Float16, Int32, UInt32, Int16, UInt16, Int8, UInt8], (f,res_idx) in [(simd_shuffle_down, 1), (simd_shuffle_up, 32)]
     function kernel(a::MtlDeviceVector{T}, b::MtlDeviceVector{T}) where T
         idx = thread_position_in_grid_1d()
         idx_in_simd = thread_index_in_simdgroup()
@@ -638,11 +638,11 @@ end
         if simd_idx == 1
             value = temp[idx_in_simd]
 
-            value = value + simd_shuffle_down(value, 16)
-            value = value + simd_shuffle_down(value,  8)
-            value = value + simd_shuffle_down(value,  4)
-            value = value + simd_shuffle_down(value,  2)
-            value = value + simd_shuffle_down(value,  1)
+            value = value + f(value, 16)
+            value = value + f(value,  8)
+            value = value + f(value,  4)
+            value = value + f(value,  2)
+            value = value + f(value,  1)
 
             b[idx] = value
         end
@@ -656,7 +656,7 @@ end
 
     rand!(a, (1:4))
     Metal.@sync @metal threads=32 kernel(dev_a, dev_b)
-    @test sum(a) ≈ b[1]
+    @test sum(a) ≈ b[res_idx]
 end
 
 @testset "matrix functions" begin