Remove obsolete MtlLargerDeviceArray (#574)

Author: christiangnrd

src/device/intrinsics/memory.jl

@@ -48,55 +48,3 @@ end
         call_function(llvm_f, Core.LLVMPtr{T,AS.ThreadGroup})
     end
 end
-
-
-## device array wrapper extending small element types
-
-struct MtlLargerDeviceArray{T,N,A} <: DenseArray{T,N}
-    x::MtlDeviceArray{UInt32,N,A}
-end
-
-Base.elsize(::Type{<:MtlLargerDeviceArray{T}}) where {T} = sizeof(UInt32)
-
-Base.size(g::MtlLargerDeviceArray) = size(g.x)
-Base.sizeof(x::MtlLargerDeviceArray) = Base.elsize(x) * length(x)
-
-Base.pointer(x::MtlLargerDeviceArray{T,<:Any,A}) where {T,A} =
-    Base.unsafe_convert(Core.LLVMPtr{T,A}, x)
-@inline function Base.pointer(x::MtlLargerDeviceArray{T,<:Any,A}, i::Integer) where {T,A}
-    Base.unsafe_convert(Core.LLVMPtr{T,A}, x) + Base._memory_offset(x, i)
-end
-
-Base.unsafe_convert(::Type{Core.LLVMPtr{T,A}}, x::MtlLargerDeviceArray{T,<:Any,A}) where {T,A} =
-    reinterpret(Core.LLVMPtr{T,A}, Base.unsafe_convert(Core.LLVMPtr{UInt32,A}, x.x))
-
-Base.@propagate_inbounds Base.getindex(A::MtlLargerDeviceArray{T}, i1::Integer) where {T} =
-    arrayref(A, i1)
-Base.@propagate_inbounds Base.setindex!(A::MtlLargerDeviceArray{T}, x, i1::Integer) where {T} =
-    arrayset(A, convert(T, x)::T, i1)
-
-# preserve the specific integer type when indexing device arrays,
-# to avoid extending 32-bit hardware indices to 64-bit.
-Base.to_index(::MtlLargerDeviceArray, i::Integer) = i
-
-# Base doesn't like Integer indices, so we need our own ND get and setindex! routines.
-# See also: https://github.com/JuliaLang/julia/pull/42289
-Base.@propagate_inbounds Base.getindex(A::MtlLargerDeviceArray,
-                                       I::Union{Integer, CartesianIndex}...) =
-    A[Base._to_linear_index(A, to_indices(A, I)...)]
-Base.@propagate_inbounds Base.setindex!(A::MtlLargerDeviceArray, x,
-                                        I::Union{Integer, CartesianIndex}...) =
-    A[Base._to_linear_index(A, to_indices(A, I)...)] = x
-
-@inline function arrayref(A::MtlLargerDeviceArray{T}, index::Integer) where {T}
-    @boundscheck checkbounds(A, index)
-    align = Base.datatype_alignment(T)
-    unsafe_load(pointer(A), index, Val(align))
-end
-
-@inline function arrayset(A::MtlLargerDeviceArray{T}, x::T, index::Integer) where {T}
-    @boundscheck checkbounds(A, index)
-    align = Base.datatype_alignment(T)
-    unsafe_store!(pointer(A), x, index, Val(align))
-    return A
-end

src/device/intrinsics/simd.jl

@@ -11,7 +11,7 @@ for (jltype, suffix) in ((:Float16, "f16"), (:Float32, "f32"))
     for as in (AS.Device, AS.ThreadGroup)
         @eval begin
             @device_function simdgroup_load(
-                data::Union{MtlDeviceArray{$jltype, <:Any, $as}, MtlLargerDeviceArray{$jltype, <:Any, $as}},
+                data::MtlDeviceArray{$jltype, <:Any, $as},
                 matrix_origin::NTuple{2, Int64} = (1, 1),
             ) = @typed_ccall($"air.simdgroup_matrix_8x8_load.v64$suffix.p$as$suffix",
                 llvmcall, NTuple{64, VecElement{$jltype}},
@@ -20,7 +20,7 @@ for (jltype, suffix) in ((:Float16, "f16"), (:Float32, "f32"))
 
             @device_function simdgroup_store(
                 src::NTuple{64, VecElement{$jltype}},
-                dest::Union{MtlDeviceArray{$jltype, <:Any, $as}, MtlLargerDeviceArray{$jltype, <:Any, $as}},
+                dest::MtlDeviceArray{$jltype, <:Any, $as},
                 matrix_origin::NTuple{2, Int64} = (1, 1),
             ) = @typed_ccall($"air.simdgroup_matrix_8x8_store.v64$suffix.p$as$suffix",
                 llvmcall, Cvoid,