This is an automated email from the ASF dual-hosted git repository.
lunderberg pushed a commit to branch main
in repository https://gitbox.apache.org/repos/asf/tvm.git
The following commit(s) were added to refs/heads/main by this push:
new 0d683284b0 [Unittest][Metal] Add minimal metal functionality test to
CI (#15756)
0d683284b0 is described below
commit 0d683284b0a5596369e2a7acf8179d02aa89d893
Author: Eric Lunderberg <[email protected]>
AuthorDate: Wed Sep 27 14:51:06 2023 -0500
[Unittest][Metal] Add minimal metal functionality test to CI (#15756)
* [Unittest][Metal] Add minimal metal functionality test to CI
Prior to this commit, the CI compiled TVM with `USE_METAL=ON` on OSX,
as defined in `conda/recipe/build.sh`, but did not validate the
execution of any generated metal kernels. As a result, breakage could
occur without being caught by the CI, such as found following
https://github.com/apache/tvm/pull/15103.
This commit adds the execution of a single metal kernel as a minimal
functionality test of the metal backend.
* CI testing, attempt a compile-only test case
* CI testing, moved intentional failure from test-case to contrib.xcode
* Move intentional failure point into codegen
* ci bump
* Removing the intentional failure during metallib compilation
---
.github/workflows/main.yml | 8 ++++++
tests/python/unittest/test_allreduce.py | 44 +++++++++++++++++++++++++++++++++
2 files changed, 52 insertions(+)
diff --git a/.github/workflows/main.yml b/.github/workflows/main.yml
index a4a30fe19a..7b4b8d826f 100644
--- a/.github/workflows/main.yml
+++ b/.github/workflows/main.yml
@@ -69,6 +69,14 @@ jobs:
shell: bash -l {0}
run: >-
python -m pytest -v tests/python/all-platform-minimal-test
+ - name: Minimal Metal Compile-Only
+ shell: bash -l {0}
+ run: >-
+ python -m pytest -v -s
'tests/python/unittest/test_allreduce.py::test_allreduce_sum_compile'
+ - name: Minimal Metal Compile-and-Run
+ shell: bash -l {0}
+ run: >-
+ python -m pytest -v -s
'tests/python/unittest/test_allreduce.py::test_allreduce_sum[dims0-metal]'
- name: Test iOS RPC
shell: bash -l {0}
run: >-
diff --git a/tests/python/unittest/test_allreduce.py
b/tests/python/unittest/test_allreduce.py
index 708384daf0..fed4e4c04d 100644
--- a/tests/python/unittest/test_allreduce.py
+++ b/tests/python/unittest/test_allreduce.py
@@ -19,6 +19,8 @@ import tvm.testing
import numpy as np
from tvm.script import tir as T
+import pytest
+
@T.prim_func
def reduce(a: T.handle, b: T.handle, d1: T.int32, d2: T.int32, d3: T.int32) ->
None:
@@ -82,6 +84,48 @@ def test_allreduce_sum(dims, target, dev):
tvm.testing.assert_allclose(b.numpy(), b_np, rtol=1e-6, atol=1e-6)
+define_metal_compile_callback = tvm.testing.parameter(True, False)
+
+
[email protected]
+def optional_metal_compile_callback(define_metal_compile_callback):
+ name = "tvm_callback_metal_compile"
+ cached = tvm.get_global_func(name, allow_missing=True)
+
+ if define_metal_compile_callback:
+
+ @tvm.register_func(name, override=True)
+ def compile_metal(src, target):
+ return tvm.contrib.xcode.compile_metal(src, sdk="macosx")
+
+ yield
+
+ if define_metal_compile_callback:
+ if cached is None:
+ tvm._ffi.registry.remove_global_func(name)
+ else:
+ tvm.register_func(name, cached, override=True)
+
+
[email protected]_metal(support_required="compile-only")
+def test_allreduce_sum_compile(optional_metal_compile_callback):
+ # Disable the parametrization over dims, at least for now
+ dims = (1, 1, 2)
+ target = "metal"
+
+ d1, d2, d3 = dims
+ _, _, _d1, _d2, _d3 = reduce.params
+ mod = reduce.specialize({_d1: d1, _d2: d2, _d3: d3})
+ sch = tvm.tir.Schedule(mod)
+ blk = sch.get_block("reduce")
+ i, j, k, l = sch.get_loops(blk)
+ sch.bind(i, "blockIdx.x")
+ sch.bind(j, "threadIdx.z")
+ sch.bind(k, "threadIdx.y")
+ sch.bind(l, "threadIdx.x")
+ tvm.build(sch.mod["main"], target=target)
+
+
@tvm.testing.parametrize_targets("cuda", "metal")
def test_allreduce_max(dims, target, dev):
d1, d2, d3 = dims
