Revert "Preserve some alignments"

This reverts commit 3ea8e1d.

ext/SpecialFunctionsExt.jl

@@ -5,7 +5,6 @@ using SpecialFunctions
 
 # math functionality corresponding to SpecialFunctions.jl
 
-# runic: off
 ## error function
 
 const tiny = 1f-30
@@ -54,7 +53,6 @@ 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/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,7 +4,6 @@
     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)),
@@ -18,7 +17,7 @@ const atomic_type_names = Dict(
     :UInt64  => "i64",
     :Float32 => "f32"
 )
-# runic: on
+
 
 ## low-level functions
 
@@ -61,7 +60,6 @@ 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],
@@ -71,7 +69,6 @@ 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]
@@ -109,7 +106,6 @@ end
 
 # copied from CUDA.jl -- should be generalized or integrated with Base
 
-# runic: off
 const inplace_ops = Dict(
     :(+=)   => :(+),
     :(-=)   => :(-),
@@ -125,7 +121,6 @@ const inplace_ops = Dict(
     :(>>=)  => :(>>),
     :(<<=)  => :(<<),
 )
-# runic: on
 
 struct AtomicError <: Exception
     msg::AbstractString
@@ -211,7 +206,6 @@ 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]),
@@ -223,7 +217,6 @@ 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,7 +86,6 @@ Returns `a * b + c`.
 
 ## SIMD Shuffle Up/Down
 
-# runic: off
 simd_shuffle_map = ((Float32, "f32"),
                     (Float16, "f16"),
                     (Int32,   "s.i32"),
@@ -95,7 +94,6 @@ 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,7 +3,6 @@ using CEnum
 
 export MemoryFlags, threadgroup_barrier, simdgroup_barrier
 
-# runic: off
 """
     MemoryFlags
 
@@ -32,7 +31,6 @@ Possible values:
     MemoryFlagTexture               = 4
     MemoryFlagThreadGroup_ImgBlock  = 8
 end
-# runic: on
 
 
 """

src/device/pointer.jl

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