Add .git-blame-ignore-revs and a few other fixes

.git-blame-ignore-revs

@@ -0,0 +1,2 @@
+# Whitespace fixes (#346)
+28881827037201fd8c8b64bd20efc7a572ab1564

docs/src/faq/contributing.md

@@ -43,6 +43,9 @@ so you should not have to define them. If using a struct for the first time in a
 interface, remember to add tests! Objective-C object methods and constructors are not yet automtically
 generatied, so any contributions there are welcome.
 
+It is also recommended to follow [these steps](https://github.com/fredrikekre/Runic.jl?tab=readme-ov-file#ignore-formatting-commits-in-git-blame) from the Runic.jl documentation
+in your local development repository so that formatting commits are ignored in blame.
+
 ## Mapping to Metal Intrinsics
 
 Some Metal functions map directly to Apple intermediate representation intrinsics. In this

ext/SpecialFunctionsExt.jl

@@ -5,6 +5,7 @@ using SpecialFunctions
 
 # math functionality corresponding to SpecialFunctions.jl
 
+# runic: off
 ## error function
 
 const tiny = 1f-30
@@ -53,6 +54,7 @@ const sb2  =  4.51839523f+01
 const sb3  =  4.72810211f+01
 const sb4  =  8.93033314f+00
 
+# runic: on
 
 # Implementation of `erf(::Float32)` from openlibm's `erfcf`
 # https://github.com/JuliaMath/openlibm/blob/12f5ffcc990e16f4120d4bf607185243f5affcb8/src/s_erff.c

src/broadcast.jl

@@ -17,7 +17,7 @@ BroadcastStyle(::MtlArrayStyle{N, S1},
     MtlArrayStyle{N, SharedStorage}()
 
 # allocation of output arrays
-Base.similar(bc::Broadcasted{MtlArrayStyle{N,S}}, ::Type{T}, dims) where {T,N,S} =
+Base.similar(::Broadcasted{MtlArrayStyle{N, S}}, ::Type{T}, dims) where {T, N, S} =
     similar(MtlArray{T,length(dims),S}, dims)
 
 # a static version of CartesianIndices that helps avoiding integer division

src/compiler/execution.jl

@@ -134,8 +134,8 @@ Adapt.adapt_structure(to::Adaptor, r::Base.RefValue) = MtlRefValue(adapt(to, r[]
 # broadcast sometimes passes a ref(type), resulting in a GPU-incompatible DataType box.
 # avoid that by using a special kind of ref that knows about the boxed type.
 struct MtlRefType{T} <: Ref{DataType} end
-Base.getindex(r::MtlRefType{T}) where T = T
-Adapt.adapt_structure(to::Adaptor, r::Base.RefValue{<:Union{DataType,Type}}) =
+Base.getindex(::MtlRefType{T}) where {T} = T
+Adapt.adapt_structure(::Adaptor, r::Base.RefValue{<:Union{DataType, Type}}) =
     MtlRefType{r[]}()
 
 # case where type is the function being broadcasted

src/compiler/library.jl

@@ -17,7 +17,7 @@ using UUIDs: UUID
 using Printf: @printf
 using CodecBzip2: Bzip2Compressor, Bzip2Decompressor, Bzip2DecompressorStream
 
-
+# runic: off
 ## enums
 
 @cenum FileType::UInt8 begin
@@ -54,7 +54,7 @@ end
     PROGRAM_INTERSECTION = 6
     PROGRAM_NONE         = 255
 end
-
+# runic: on
 
 ## structures
 

src/device/intrinsics/atomics.jl

@@ -4,6 +4,7 @@
     memory_order_relaxed = 0
 end
 
+# runic: off
 # XXX: the integers should come from some enum
 const atomic_memory_names = Dict(
     AS.Device      => ("global", Int32(2)),
@@ -17,7 +18,7 @@ const atomic_type_names = Dict(
     :UInt64  => "i64",
     :Float32 => "f32"
 )
-
+# runic: on
 
 ## low-level functions
 
@@ -60,6 +61,7 @@ for typ in (:Int32, :UInt32, :Float32), as in (AS.Device, AS.ThreadGroup)
     end
 end
 
+# runic: off
 const atomic_fetch_and_modify = [
     :add => [:Int32, :UInt32, :Float32],
     :sub => [:Int32, :UInt32, :Float32],
@@ -69,6 +71,7 @@ const atomic_fetch_and_modify = [
     :or  => [:Int32, :UInt32],
     :xor => [:Int32, :UInt32]
 ]
+# runic: on
 
 for (op, types) in atomic_fetch_and_modify, typ in types, as in (AS.Device, AS.ThreadGroup)
     typnam = atomic_type_names[typ]
@@ -106,6 +109,7 @@ end
 
 # copied from CUDA.jl -- should be generalized or integrated with Base
 
+# runic: off
 const inplace_ops = Dict(
     :(+=)   => :(+),
     :(-=)   => :(-),
@@ -121,6 +125,7 @@ const inplace_ops = Dict(
     :(>>=)  => :(>>),
     :(<<=)  => :(<<),
 )
+# runic: on
 
 struct AtomicError <: Exception
     msg::AbstractString
@@ -206,6 +211,7 @@ end
 @inline atomic_arrayref(A::AbstractArray, I::Integer) = atomic_load_explicit(pointer(A, I))
 @inline atomic_arrayset(A::AbstractArray{T}, I::Integer, ::Nothing, val) where T =
     atomic_store_explicit(pointer(A, I), convert(T, val))
+# runic: off
 for (op,impl,typ) in [(:(+), :(atomic_fetch_add_explicit), [:UInt32,:Int32,:Float32]),
                       (:(-), :(atomic_fetch_sub_explicit), [:UInt32,:Int32,:Float32]),
                       (:(&), :(atomic_fetch_and_explicit), [:UInt32,:Int32]),
@@ -217,6 +223,7 @@ for (op,impl,typ) in [(:(+), :(atomic_fetch_add_explicit), [:UInt32,:Int32,:Floa
                                   val::T) where {T<:Union{$(typ...)}} =
         $impl(pointer(A, I), val)
 end
+# runic: on
 
 # native atomics that are not supported on all devices
 @inline function atomic_arrayset(A::AbstractArray{T}, I::Integer, op::typeof(+),

src/device/intrinsics/math.jl

@@ -333,7 +333,7 @@ end
 
 
 ### Integer Intrinsics
-
+# runic: off
 @device_override Base.abs(x::Int64)   = ccall("extern air.abs.s.i64", llvmcall, Int64, (Int64,), x)
 @device_override Base.abs(x::UInt64)  = ccall("extern air.abs.u.i64", llvmcall, UInt64, (UInt64,), x)
 @device_override Base.abs(x::Int32)   = ccall("extern air.abs.s.i32", llvmcall, Int32, (Int32,), x)
@@ -396,7 +396,7 @@ end
 @device_function reverse_bits(x::UInt16)  = ccall("extern air.reverse_bits.i16", llvmcall, UInt16, (UInt16,), x)
 @device_function reverse_bits(x::Int8)    = ccall("extern air.reverse_bits.i8", llvmcall, Int8, (Int8,), x)
 @device_function reverse_bits(x::UInt8)   = ccall("extern air.reverse_bits.i8", llvmcall, UInt8, (UInt8,), x)
-
+# runic: on
 
 function _mulhi(a::Int64, b::Int64)
     shift = sizeof(a) * 4

src/device/intrinsics/simd.jl

@@ -86,6 +86,7 @@ Returns `a * b + c`.
 
 ## SIMD Shuffle Up/Down
 
+# runic: off
 simd_shuffle_map = ((Float32, "f32"),
                     (Float16, "f16"),
                     (Int32,   "s.i32"),
@@ -94,6 +95,7 @@ simd_shuffle_map = ((Float32, "f32"),
                     (UInt16,  "u.i16"),
                     (Int8,    "s.i8"),
                     (UInt8,   "u.i8"))
+# runic: on
 
 for (jltype, suffix) in simd_shuffle_map
     @eval begin

src/device/intrinsics/synchronization.jl

@@ -3,6 +3,7 @@ using CEnum
 
 export MemoryFlags, threadgroup_barrier, simdgroup_barrier
 
+# runic: off
 """
     MemoryFlags
 
@@ -31,6 +32,7 @@ Possible values:
     MemoryFlagTexture               = 4
     MemoryFlagThreadGroup_ImgBlock  = 8
 end
+# runic: on
 
 
 """

src/device/pointer.jl

@@ -1,3 +1,4 @@
+# runic: off
 module AS
     const Generic               = 0 # No Generic address space?
     const Device                = 1 # Checked

src/version.jl

@@ -86,7 +86,7 @@ function metal_support()
         v"3.0"
     elseif macos >= v"12"
         v"2.4"
-    elseif macos v> v"11"
+    elseif macos >= v"11"
         v"2.3"
     elseif macos >= v"10.15"
         v"2.2"