From 65bd048699305b4da56db7a474654751b36e6324 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Fri, 15 Sep 2023 08:38:17 -0500 Subject: [PATCH 1/6] [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. --- .github/workflows/main.yml | 4 ++++ tests/python/unittest/test_allreduce.py | 3 +++ 2 files changed, 7 insertions(+) diff --git a/.github/workflows/main.yml b/.github/workflows/main.yml index a4a30fe19aba..c0a9fec1498a 100644 --- a/.github/workflows/main.yml +++ b/.github/workflows/main.yml @@ -69,6 +69,10 @@ jobs: shell: bash -l {0} run: >- python -m pytest -v tests/python/all-platform-minimal-test + - name: Minimal Metal Functionality + 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 708384daf027..62681e16a551 100644 --- a/tests/python/unittest/test_allreduce.py +++ b/tests/python/unittest/test_allreduce.py @@ -59,6 +59,9 @@ def generate_param_sets(): @tvm.testing.parametrize_targets("cuda", "metal") def test_allreduce_sum(dims, target, dev): + print(f"Executing on target {target}") + assert target != "metal", "Deliberate failure to check that CI is executing the test" + d1, d2, d3 = dims _, _, _d1, _d2, _d3 = reduce.params mod = reduce.specialize({_d1: d1, _d2: d2, _d3: d3}) From 6f396e266ceda3ce41570e034e43f7742f9ddf98 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Wed, 20 Sep 2023 09:43:34 -0500 Subject: [PATCH 2/6] CI testing, attempt a compile-only test case --- .github/workflows/main.yml | 6 +++- tests/python/unittest/test_allreduce.py | 46 +++++++++++++++++++++++++ 2 files changed, 51 insertions(+), 1 deletion(-) diff --git a/.github/workflows/main.yml b/.github/workflows/main.yml index c0a9fec1498a..7b4b8d826f36 100644 --- a/.github/workflows/main.yml +++ b/.github/workflows/main.yml @@ -69,7 +69,11 @@ jobs: shell: bash -l {0} run: >- python -m pytest -v tests/python/all-platform-minimal-test - - name: Minimal Metal Functionality + - 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]' diff --git a/tests/python/unittest/test_allreduce.py b/tests/python/unittest/test_allreduce.py index 62681e16a551..15e2517d1f06 100644 --- a/tests/python/unittest/test_allreduce.py +++ b/tests/python/unittest/test_allreduce.py @@ -19,6 +19,8 @@ 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: @@ -85,6 +87,50 @@ 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) + + +@pytest.fixture +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) + + +@tvm.testing.requires_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" + print(f"Executing on target {target}") + assert target != "metal", "Deliberate failure to check that CI is executing the test" + + 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 From 5ee501121fcaada13f73950e1eff926fb8661cf7 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Thu, 21 Sep 2023 09:40:46 -0500 Subject: [PATCH 3/6] CI testing, moved intentional failure from test-case to contrib.xcode --- python/tvm/contrib/xcode.py | 3 +++ tests/python/unittest/test_allreduce.py | 2 -- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/python/tvm/contrib/xcode.py b/python/tvm/contrib/xcode.py index 2b68600197e4..1c17e2d9b3bb 100644 --- a/python/tvm/contrib/xcode.py +++ b/python/tvm/contrib/xcode.py @@ -165,6 +165,9 @@ def compile_metal(code, path_target=None, sdk="macosx", min_os_version=None): libbin = None else: libbin = bytearray(open(file_target, "rb").read()) + + + assert False, "Deliberate failure after metal compilation, to ensure the CI reached this step" return libbin diff --git a/tests/python/unittest/test_allreduce.py b/tests/python/unittest/test_allreduce.py index 15e2517d1f06..0353dd28e232 100644 --- a/tests/python/unittest/test_allreduce.py +++ b/tests/python/unittest/test_allreduce.py @@ -115,8 +115,6 @@ def test_allreduce_sum_compile(optional_metal_compile_callback): # Disable the parametrization over dims, at least for now dims = (1, 1, 2) target = "metal" - print(f"Executing on target {target}") - assert target != "metal", "Deliberate failure to check that CI is executing the test" d1, d2, d3 = dims _, _, _d1, _d2, _d3 = reduce.params From 84ff31c58f8463cac9d5287be581b6dbbead3836 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Thu, 21 Sep 2023 14:47:41 -0500 Subject: [PATCH 4/6] Move intentional failure point into codegen --- python/tvm/contrib/xcode.py | 3 --- src/target/source/codegen_metal.cc | 5 +++++ 2 files changed, 5 insertions(+), 3 deletions(-) diff --git a/python/tvm/contrib/xcode.py b/python/tvm/contrib/xcode.py index 1c17e2d9b3bb..2b68600197e4 100644 --- a/python/tvm/contrib/xcode.py +++ b/python/tvm/contrib/xcode.py @@ -165,9 +165,6 @@ def compile_metal(code, path_target=None, sdk="macosx", min_os_version=None): libbin = None else: libbin = bytearray(open(file_target, "rb").read()) - - - assert False, "Deliberate failure after metal compilation, to ensure the CI reached this step" return libbin diff --git a/src/target/source/codegen_metal.cc b/src/target/source/codegen_metal.cc index b8c30691e21f..848fa3981078 100644 --- a/src/target/source/codegen_metal.cc +++ b/src/target/source/codegen_metal.cc @@ -172,6 +172,11 @@ void CodeGenMetal::BindThreadIndex(const IterVar& iv) { } void CodeGenMetal::PrintType(DataType t, std::ostream& os) { // NOLINT(*) + os << "\n" + << "Deliberate insertion of code " + << "that should flagrantly fail to compile, " + << "and whose failure should be caught in CI" + << "\n"; int lanes = t.lanes(); if (t.is_handle()) { ICHECK_EQ(lanes, 1) << "do not yet support vector types"; From 3e8684786f173be2d2a1fed31593a1e6221b588f Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Sat, 23 Sep 2023 07:42:44 -0500 Subject: [PATCH 5/6] ci bump From 1725292ac201fa56ddd9278086e39e480684f502 Mon Sep 17 00:00:00 2001 From: Eric Lunderberg Date: Mon, 25 Sep 2023 10:17:53 -0500 Subject: [PATCH 6/6] Removing the intentional failure during metallib compilation --- src/target/source/codegen_metal.cc | 5 ----- tests/python/unittest/test_allreduce.py | 3 --- 2 files changed, 8 deletions(-) diff --git a/src/target/source/codegen_metal.cc b/src/target/source/codegen_metal.cc index 848fa3981078..b8c30691e21f 100644 --- a/src/target/source/codegen_metal.cc +++ b/src/target/source/codegen_metal.cc @@ -172,11 +172,6 @@ void CodeGenMetal::BindThreadIndex(const IterVar& iv) { } void CodeGenMetal::PrintType(DataType t, std::ostream& os) { // NOLINT(*) - os << "\n" - << "Deliberate insertion of code " - << "that should flagrantly fail to compile, " - << "and whose failure should be caught in CI" - << "\n"; int lanes = t.lanes(); if (t.is_handle()) { ICHECK_EQ(lanes, 1) << "do not yet support vector types"; diff --git a/tests/python/unittest/test_allreduce.py b/tests/python/unittest/test_allreduce.py index 0353dd28e232..fed4e4c04d32 100644 --- a/tests/python/unittest/test_allreduce.py +++ b/tests/python/unittest/test_allreduce.py @@ -61,9 +61,6 @@ def generate_param_sets(): @tvm.testing.parametrize_targets("cuda", "metal") def test_allreduce_sum(dims, target, dev): - print(f"Executing on target {target}") - assert target != "metal", "Deliberate failure to check that CI is executing the test" - d1, d2, d3 = dims _, _, _d1, _d2, _d3 = reduce.params mod = reduce.specialize({_d1: d1, _d2: d2, _d3: d3})