wehub-resource-sync
463 lines
19 KiB
Python
463 lines
19 KiB
Python
# Licensed to the Apache Software Foundation (ASF) under one
|
|
# or more contributor license agreements. See the NOTICE file
|
|
# distributed with this work for additional information
|
|
# regarding copyright ownership. The ASF licenses this file
|
|
# to you under the Apache License, Version 2.0 (the
|
|
# "License"); you may not use this file except in compliance
|
|
# with the License. You may obtain a copy of the License at
|
|
#
|
|
# http://www.apache.org/licenses/LICENSE-2.0
|
|
#
|
|
# Unless required by applicable law or agreed to in writing,
|
|
# software distributed under the License is distributed on an
|
|
# "AS IS" BASIS, WITHOUT WARRANTIES OR CONDITIONS OF ANY
|
|
# KIND, either express or implied. See the License for the
|
|
# specific language governing permissions and limitations
|
|
# under the License.
|
|
# pylint: disable=missing-function-docstring,missing-module-docstring
|
|
# ruff: noqa: F401, F841
|
|
import pytest
|
|
|
|
import tvm
|
|
import tvm.testing
|
|
from tvm import te, tirx
|
|
from tvm.s_tir.schedule.testing import (
|
|
assert_structural_equal_ignore_global_symbol,
|
|
verify_trace_roundtrip,
|
|
)
|
|
from tvm.script import tirx as T
|
|
from tvm.tirx.expr import IntImm
|
|
|
|
# pylint: disable=no-member,invalid-name,unused-variable
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, (128, 128, 128))
|
|
B = T.match_buffer(b, (128, 128, 128))
|
|
for i, j, k in T.grid(128, 128, 128):
|
|
with T.sblock("B"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i, j, k])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * 2.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_symbolic(a: T.handle, b: T.handle, n: T.int32) -> None:
|
|
A = T.match_buffer(a, (128, 128, n))
|
|
B = T.match_buffer(b, (128, 128, n))
|
|
for i, j, k in T.grid(128, 128, n):
|
|
with T.sblock("B"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i, j, k])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * 2.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_with_anno(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, (128, 128, 128))
|
|
B = T.match_buffer(b, (128, 128, 128))
|
|
for i, j in T.grid(128, 128):
|
|
for k in T.serial(0, 128, annotations={"useless_annotation": True}):
|
|
with T.sblock("B"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i, j, k])
|
|
T.reads([A[vi, vj, vk]])
|
|
T.writes([B[vi, vj, vk]])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * 2.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_with_thread_binding(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, (128, 128, 128))
|
|
B = T.match_buffer(b, (128, 128, 128))
|
|
for i, j in T.grid(128, 128):
|
|
for k in T.thread_binding(0, 128, thread="threadIdx.x"):
|
|
with T.sblock("B"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i, j, k])
|
|
T.reads([A[vi, vj, vk]])
|
|
T.writes([B[vi, vj, vk]])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * 2.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_with_opaque_block(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, (128, 128, 128))
|
|
B = T.match_buffer(b, (128, 128, 128))
|
|
for i, j, k in T.grid(128, 128, 128):
|
|
with T.sblock("opaque"):
|
|
T.reads([A[i, j, k]])
|
|
T.writes([B[i, j, k]])
|
|
with T.sblock("B"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i, j, k])
|
|
T.reads([A[vi, vj, vk]])
|
|
T.writes([B[vi, vj, vk]])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * 2.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_partition_with_opaque_block(a: T.handle, b: T.handle) -> None:
|
|
B = T.match_buffer(b, [128, 128, 128])
|
|
A = T.match_buffer(a, [128, 128, 128])
|
|
with T.sblock("root"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("opaque_i_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("opaque_i0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i0, j, k in T.grid(112, 128, 128):
|
|
with T.sblock("opaque_i0"):
|
|
T.reads(A[i0, j, k])
|
|
T.writes(B[i0, j, k])
|
|
with T.sblock("B_i0"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i0, j, k])
|
|
T.reads(A[0:112, 0:128, 0:128])
|
|
T.writes(B[0:112, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("opaque_i1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i1 in range(112, 128):
|
|
for j, k in T.grid(128, 128):
|
|
with T.sblock("opaque_i1"):
|
|
T.reads(A[i1, j, k])
|
|
T.writes(B[i1, j, k])
|
|
with T.sblock("B_i1"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i1, j, k])
|
|
T.reads(A[112:128, 0:128, 0:128])
|
|
T.writes(B[112:128, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_loop_partition_case0(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, [128, 128, 128])
|
|
B = T.match_buffer(b, [128, 128, 128])
|
|
with T.sblock("root"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_i_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_i0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i0 in range(2):
|
|
with T.sblock("B_i0_j_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_i0_j0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j0, k in T.grid(4, 128):
|
|
with T.sblock("B_i0_j0"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i0, j0, k])
|
|
T.reads(A[0:2, 0:4, 0:128])
|
|
T.writes(B[0:2, 0:4, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i0_j1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j1 in range(4, 36):
|
|
for k in range(128):
|
|
with T.sblock("B_i0_j1"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i0, j1, k])
|
|
T.reads(A[0:2, 4:36, 0:128])
|
|
T.writes(B[0:2, 4:36, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i0_j2_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j2 in range(36, 128):
|
|
for k in range(128):
|
|
with T.sblock("B_i0_j2"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i0, j2, k])
|
|
T.reads(A[0:2, 36:128, 0:128])
|
|
T.writes(B[0:2, 36:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i1 in range(2, 3):
|
|
for j, k in T.grid(128, 128):
|
|
with T.sblock("B_i1"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i1, j, k])
|
|
T.reads(A[2, 0:128, 0:128])
|
|
T.writes(B[2, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i2_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i2 in range(3, 67):
|
|
for j, k in T.grid(128, 128):
|
|
with T.sblock("B_i2"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i2, j, k])
|
|
T.reads(A[3:67, 0:128, 0:128])
|
|
T.writes(B[3:67, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i3_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i3 in range(67, 128):
|
|
for j, k in T.grid(128, 128):
|
|
with T.sblock("B_i3"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i3, j, k])
|
|
T.reads(A[67:128, 0:128, 0:128])
|
|
T.writes(B[67:128, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def elementwise_loop_partition_case1(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, [128, 128, 128])
|
|
B = T.match_buffer(b, [128, 128, 128])
|
|
with T.sblock("root"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_i_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_i0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i0, j, k in T.grid(63, 128, 128):
|
|
with T.sblock("B_i0"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i0, j, k])
|
|
T.reads(A[0:63, 0:128, 0:128])
|
|
T.writes(B[0:63, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i1 in range(63, 64):
|
|
for j in range(128):
|
|
with T.sblock("B_i1_k_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_i1_k0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for k0 in range(1):
|
|
with T.sblock("B_i1_k0"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i1, j, k0])
|
|
T.reads(A[63, 0:128, 0])
|
|
T.writes(B[63, 0:128, 0])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i1_k1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for k1 in range(1, 65):
|
|
with T.sblock("B_i1_k1"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i1, j, k1])
|
|
T.reads(A[63, 0:128, 1:65])
|
|
T.writes(B[63, 0:128, 1:65])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i1_k2_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for k2 in range(65, 128):
|
|
with T.sblock("B_i1_k2"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i1, j, k2])
|
|
T.reads(A[63, 0:128, 65:128])
|
|
T.writes(B[63, 0:128, 65:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
with T.sblock("B_i2_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for i2 in range(64, 128):
|
|
for j, k in T.grid(128, 128):
|
|
with T.sblock("B_i2"):
|
|
vi, vj, vk = T.axis.remap("SSS", [i2, j, k])
|
|
T.reads(A[64:128, 0:128, 0:128])
|
|
T.writes(B[64:128, 0:128, 0:128])
|
|
B[vi, vj, vk] = A[vi, vj, vk] * T.float32(2)
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def opaque_access(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, [16, 16], "float32")
|
|
B = T.match_buffer(b, [16, 16], "float32")
|
|
for i, j in T.grid(16, 16):
|
|
with T.sblock("A"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
T.reads([])
|
|
T.writes([A[0:16, 0:16]])
|
|
A[vi, vj] = 1
|
|
for i, j in T.grid(16, 16):
|
|
with T.sblock("B"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
T.reads([])
|
|
T.writes([B[0:16, 0:16]])
|
|
T.evaluate(T.tvm_fill_fragment(B.data, 16, 16, 16, 0, vi * 16 + vj, dtype="handle"))
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def opaque_access_loop_partition(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, (16, 16))
|
|
B = T.match_buffer(b, (16, 16))
|
|
for i in range(16):
|
|
with T.sblock("A_j_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("A_j0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j0 in range(12):
|
|
with T.sblock("A_j0"):
|
|
vi, vj = T.axis.remap("SS", [i, j0])
|
|
T.reads()
|
|
T.writes(A[0:16, 0:12])
|
|
A[vi, vj] = T.float32(1)
|
|
with T.sblock("A_j1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j1 in range(12, 16):
|
|
with T.sblock("A_j1"):
|
|
vi, vj = T.axis.remap("SS", [i, j1])
|
|
T.reads()
|
|
T.writes(A[0:16, 12:16])
|
|
A[vi, vj] = T.float32(1)
|
|
for i in range(16):
|
|
with T.sblock("B_j_common"):
|
|
T.reads()
|
|
T.writes()
|
|
with T.sblock("B_j0_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j0 in range(12):
|
|
with T.sblock("B_j0"):
|
|
vi, vj = T.axis.remap("SS", [i, j0])
|
|
T.reads()
|
|
T.writes(B[0:16, 0:16])
|
|
T.tvm_fill_fragment(B.data, 16, 16, 16, 0, vi * 16 + vj)
|
|
with T.sblock("B_j1_partition"):
|
|
T.reads()
|
|
T.writes()
|
|
for j1 in range(12, 16):
|
|
with T.sblock("B_j1"):
|
|
vi, vj = T.axis.remap("SS", [i, j1])
|
|
T.reads()
|
|
T.writes(B[0:16, 0:16])
|
|
T.tvm_fill_fragment(B.data, 16, 16, 16, 0, vi * 16 + vj)
|
|
|
|
|
|
# pylint: enable=no-member,invalid-name,unused-variable
|
|
|
|
|
|
def test_loop_partition():
|
|
sch = tvm.s_tir.Schedule(elementwise, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
i, j, k = sch.get_loops(block_b)
|
|
sch.loop_partition(i, factors=[2, 1, 64])
|
|
|
|
block_b_partition = sch.get_sblock("B_i0")
|
|
i, j, k = sch.get_loops(block_b_partition)
|
|
loops = sch.loop_partition(j, factors=[4, 32])
|
|
|
|
assert_structural_equal_ignore_global_symbol(elementwise_loop_partition_case0, sch.mod["main"])
|
|
verify_trace_roundtrip(sch=sch, mod=elementwise)
|
|
|
|
|
|
def test_partition_with_inferred_factor():
|
|
sch = tvm.s_tir.Schedule(elementwise, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
i, j, k = sch.get_loops(block_b)
|
|
sch.loop_partition(i, factors=[None, 1, 64])
|
|
|
|
block_b_partition = sch.get_sblock("B_i1")
|
|
i, j, k = sch.get_loops(block_b_partition)
|
|
sch.loop_partition(k, factors=[1, 64, None])
|
|
|
|
assert_structural_equal_ignore_global_symbol(elementwise_loop_partition_case1, sch.mod["main"])
|
|
verify_trace_roundtrip(sch=sch, mod=elementwise)
|
|
|
|
|
|
def test_partition_with_opaque_block():
|
|
sch = tvm.s_tir.Schedule(elementwise_with_opaque_block, debug_mask="all")
|
|
block_opaque = sch.get_sblock("opaque")
|
|
i, _, _ = sch.get_loops(block_opaque)
|
|
sch.loop_partition(i, factors=[None, 16])
|
|
assert_structural_equal_ignore_global_symbol(
|
|
elementwise_partition_with_opaque_block, sch.mod["main"]
|
|
)
|
|
verify_trace_roundtrip(sch=sch, mod=elementwise_with_opaque_block)
|
|
|
|
|
|
def test_partition_with_opaque_access():
|
|
sch = tvm.s_tir.Schedule(opaque_access, debug_mask="all")
|
|
block_a = sch.get_sblock("A")
|
|
_, j = sch.get_loops(block_a)
|
|
sch.loop_partition(j, factors=[None, 4])
|
|
block_b = sch.get_sblock("B")
|
|
_, j = sch.get_loops(block_b)
|
|
sch.loop_partition(j, factors=[None, 4])
|
|
assert_structural_equal_ignore_global_symbol(opaque_access_loop_partition, sch.mod["main"])
|
|
verify_trace_roundtrip(sch=sch, mod=opaque_access)
|
|
|
|
|
|
def test_partition_int64_extent_with_mixed_factors():
|
|
def _create_prim_func():
|
|
m = te.const(384, "int64")
|
|
A = te.placeholder((m,), name="A", dtype="float32")
|
|
B = te.compute((m,), lambda i: A[i] + 1, name="B")
|
|
return te.create_prim_func([A, B])
|
|
|
|
mod = _create_prim_func()
|
|
sch = tvm.s_tir.Schedule(mod, debug_mask="all")
|
|
(i,) = sch.get_loops(sch.get_sblock("B"))
|
|
sch.loop_partition(
|
|
i,
|
|
factors=[
|
|
te.const(1, "int64"),
|
|
te.const(51, "int32"),
|
|
],
|
|
)
|
|
|
|
|
|
def test_partition_fail_symbolic():
|
|
sch = tvm.s_tir.Schedule(elementwise_symbolic, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
_, _, k = sch.get_loops(block_b)
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(k, factors=[10, None])
|
|
|
|
|
|
def test_partition_fail_out_of_bound():
|
|
sch = tvm.s_tir.Schedule(elementwise, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
i, j, k = sch.get_loops(block_b)
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(i, factors=[1000, 2, 3])
|
|
|
|
|
|
def test_partition_with_non_positive_factors():
|
|
sch = tvm.s_tir.Schedule(elementwise, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
i, j, k = sch.get_loops(block_b)
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(i, factors=[-2, -64])
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(j, factors=[0, None])
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(k, factors=[None, -16])
|
|
|
|
|
|
def test_partition_fail_with_annotation():
|
|
sch = tvm.s_tir.Schedule(elementwise_with_anno, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
_, j, k = sch.get_loops(block_b)
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(k, factors=[None, 10])
|
|
|
|
|
|
def test_partition_fail_with_thread_binding():
|
|
sch = tvm.s_tir.Schedule(elementwise_with_thread_binding, debug_mask="all")
|
|
block_b = sch.get_sblock("B")
|
|
_, j, k = sch.get_loops(block_b)
|
|
with pytest.raises(tvm.s_tir.ScheduleError):
|
|
sch.loop_partition(k, factors=[None, 10])
|
|
|
|
|
|
if __name__ == "__main__":
|
|
tvm.testing.main()
|