From 793d92d1fe15e0dda0f5c8484dd9157f9fbf40d2 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Fri, 7 Jun 2024 17:08:18 +0200 Subject: [PATCH 1/5] Bump compat. --- Project.toml | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/Project.toml b/Project.toml index 8ac1ddbb0..f50b9ecce 100644 --- a/Project.toml +++ b/Project.toml @@ -37,8 +37,8 @@ ExprTools = "0.1" GPUArrays = "10.1" GPUCompiler = "0.26" KernelAbstractions = "0.9.1" -LLVM = "6, 7" -LLVMDowngrader_jll = "0.2" +LLVM = "7.2" +LLVMDowngrader_jll = "0.3" ObjectFile = "0.4" ObjectiveC = "2.1, 3" Preferences = "1" From 886567c742cb9b7a222b15c2ab8a9cf176e203d1 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Fri, 7 Jun 2024 17:08:34 +0200 Subject: [PATCH 2/5] Emit metadata to reconstruct typed pointers. --- src/compiler/compilation.jl | 41 +++++++++++++++++++++++++++++++++++++ 1 file changed, 41 insertions(+) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index c3d6ce428..3b7e1329f 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -9,6 +9,47 @@ GPUCompiler.runtime_module(::MetalCompilerJob) = Metal GPUCompiler.method_table(::MetalCompilerJob) = method_table +function GPUCompiler.finish_ir!(@nospecialize(job::MetalCompilerJob), + mod::LLVM.Module, entry::LLVM.Function) + entry = invoke(GPUCompiler.finish_ir!, + Tuple{CompilerJob{MetalCompilerTarget}, LLVM.Module, LLVM.Function}, + job, mod, entry) + + # pointer type information for typed intrinsics + # (this is consumed by the LLVM IR downgrader) + for (jltyp, llvmtyp) in (Int32 => :i32, Int64 => :i64, + Float16 => :f16, Float32 => :f32), + (as, asname) in (AS.Device => "global", AS.ThreadGroup => "local") + + # map of intrinsics to pointer operand indices and eltypes + intrinsics = Dict() + ## simd + intrinsics["simdgroup_matrix_8x8_load.v64$llvmtyp.p$as$llvmtyp"] = (1 => jltyp,) + intrinsics["simdgroup_matrix_8x8_store.v64$llvmtyp.p$as$llvmtyp"] = (2 => jltyp,) + ## atomics + for op in [:store, :load, :xchg, :add, :sub, :min, :max, :and, :or, :xor] + intrinsics["atomic.$asname.$op.$llvmtyp"] = (1 => jltyp,) + end + intrinsics["atomic.$asname.cmpxchg.weak.$llvmtyp"] = (1 => jltyp, 2 => jltyp) + + # apply metadata to the function declarations + for (intr, args) in intrinsics + fn = "air.$intr" + haskey(functions(mod), fn) || continue + f = functions(mod)[fn] + mds = [] + for (idx, typ) in args + push!(mds, ConstantInt(Int32(idx-1))) + push!(mds, null(convert(LLVMType, typ))) + end + metadata(f)["arg_eltypes"] = MDNode(mds) + end + end + + return entry +end + + ## compiler implementation (cache, configure, compile, and link) # cache of compilation caches, per device From f75127e0b9c61f5800af2c974467dcb6db0a455c Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Fri, 7 Jun 2024 17:08:45 +0200 Subject: [PATCH 3/5] Support opaque pointers. --- src/compiler/compilation.jl | 3 --- 1 file changed, 3 deletions(-) diff --git a/src/compiler/compilation.jl b/src/compiler/compilation.jl index 3b7e1329f..0bb81c372 100644 --- a/src/compiler/compilation.jl +++ b/src/compiler/compilation.jl @@ -111,9 +111,6 @@ function compile(@nospecialize(job::CompilerJob)) log = Pipe() cmd = `$(LLVMDowngrader_jll.llvm_as()) --bitcode-version=5.0 -o -` - if LLVM.version() >= v"16" - cmd = `$cmd --opaque-pointers=0` - end proc = run(pipeline(cmd, stdout=output, stderr=log, stdin=input); wait=false) close(output.in) close(log.in) From e87f07ee51fa9191e6f4d7fd06bb68fb92ad1e49 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Fri, 7 Jun 2024 17:08:54 +0200 Subject: [PATCH 4/5] Don't put types in the AST. --- src/device/intrinsics/simd.jl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/src/device/intrinsics/simd.jl b/src/device/intrinsics/simd.jl index 95bf3ae5f..79250e330 100644 --- a/src/device/intrinsics/simd.jl +++ b/src/device/intrinsics/simd.jl @@ -7,7 +7,7 @@ function convert_origin(origin::NTuple{2, Int64}) return (VecElement{Int64}(origin[1]-1), VecElement{Int64}(origin[2]-1)) end -for (jltype, suffix) in ((Float16, "f16"), (Float32, "f32")) +for (jltype, suffix) in ((:Float16, "f16"), (:Float32, "f32")) for as in (AS.Device, AS.ThreadGroup) @eval begin @device_function simdgroup_load( From 1f0bcccf3de13a1f5a22e4b57d05e3247240da53 Mon Sep 17 00:00:00 2001 From: Tim Besard Date: Sat, 8 Jun 2024 09:11:03 +0200 Subject: [PATCH 5/5] Add opaque pointer CI job. --- .buildkite/pipeline.yml | 38 ++++++++++++++++++++++++++------------ 1 file changed, 26 insertions(+), 12 deletions(-) diff --git a/.buildkite/pipeline.yml b/.buildkite/pipeline.yml index dcde7ac20..5315f3ff1 100644 --- a/.buildkite/pipeline.yml +++ b/.buildkite/pipeline.yml @@ -34,11 +34,11 @@ steps: julia: "nightly" soft_fail: true - # Test Storage modes Shared and Managed - - group: ":floppy_disk: Storage mode" - key: "storage" + # special tests + - group: ":eyes: Special" + depends_on: "julia" steps: - - label: "MtlArray with {{matrix.storage}} storage" + - label: "{{matrix.storage}} array storage" plugins: - JuliaCI/julia#v1: version: "1.10" @@ -48,7 +48,7 @@ steps: queue: "juliaecosystem" os: "macos" arch: "aarch64" - if: build.message !~ /\[skip tests\]/ && build.message !~ /\[skip storage\]/ && !build.pull_request.draft + if: build.message !~ /\[skip tests\]/ && !build.pull_request.draft timeout_in_minutes: 60 matrix: setup: @@ -57,13 +57,7 @@ steps: - "Managed" commands: | echo -e "[Metal]\ndefault_storage = \"{{matrix.storage}}\"" >LocalPreferences.toml - - - # special tests - - group: ":eyes: Special" - depends_on: "julia" - steps: - - label: "Metal API Validation" + - label: "API validation" plugins: - JuliaCI/julia#v1: version: 1.8 @@ -85,3 +79,23 @@ steps: arch: "aarch64" if: build.message !~ /\[skip tests\]/ && !build.pull_request.draft timeout_in_minutes: 60 + - label: "Opaque pointers" + plugins: + - JuliaCI/julia#v1: + version: "1.10" + - JuliaCI/julia-test#v1: + test_args: "--quickfail" + - JuliaCI/julia-coverage#v1: + codecov: true + dirs: + - src + - lib + - examples + env: + JULIA_LLVM_ARGS: '--opaque-pointers' + agents: + queue: "juliaecosystem" + os: "macos" + arch: "aarch64" + if: build.message !~ /\[skip tests\]/ && !build.pull_request.draft + timeout_in_minutes: 60