wehub-resource-sync
366 lines
12 KiB
Python
366 lines
12 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
|
|
from tvm.s_tir.schedule.testing import assert_structural_equal_ignore_global_symbol
|
|
from tvm.script import tirx as T
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def matmul(a: T.handle, b: T.handle, c: T.handle, n: T.int32) -> None:
|
|
m = T.int32()
|
|
A = T.match_buffer(a, [m, n])
|
|
B = T.match_buffer(b, [m, n])
|
|
C = T.match_buffer(c, [m, m])
|
|
|
|
for i, j, k in T.grid(m, m, n):
|
|
with T.sblock("update"):
|
|
vi, vj, vk = T.axis.remap("SSR", [i, j, k])
|
|
with T.init():
|
|
C[vi, vj] = 0.0
|
|
C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vj, vk]
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def matmul_128(a: T.handle, b: T.handle, c: T.handle) -> None:
|
|
A = T.match_buffer(a, [128, 128])
|
|
B = T.match_buffer(b, [128, 128])
|
|
C = T.match_buffer(c, [128, 128])
|
|
|
|
for i, j, k in T.grid(128, 128, 128):
|
|
with T.sblock("update"):
|
|
vi, vj, vk = T.axis.remap("SSR", [i, j, k])
|
|
with T.init():
|
|
C[vi, vj] = 0.0
|
|
C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vj, vk]
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def matmul_m_128(a: T.handle, b: T.handle, c: T.handle) -> None:
|
|
m = T.int32()
|
|
A = T.match_buffer(a, [m, 128])
|
|
B = T.match_buffer(b, [m, 128])
|
|
C = T.match_buffer(c, [m, m])
|
|
|
|
for i, j, k in T.grid(m, m, 128):
|
|
with T.sblock("update"):
|
|
vi, vj, vk = T.axis.remap("SSR", [i, j, k])
|
|
with T.init():
|
|
C[vi, vj] = 0.0
|
|
C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vj, vk]
|
|
|
|
|
|
# x is considered undefined because it appears as part of x*8,
|
|
# but not on its own
|
|
@T.prim_func(check_well_formed=False, s_tir=True)
|
|
def matmul_m_8x(a: T.handle, b: T.handle, c: T.handle) -> None:
|
|
x = T.int32()
|
|
m = T.int32()
|
|
A = T.match_buffer(a, [m, x * 8])
|
|
B = T.match_buffer(b, [m, x * 8])
|
|
C = T.match_buffer(c, [m, m])
|
|
|
|
for i, j, k in T.grid(m, m, x * 8):
|
|
with T.sblock("update"):
|
|
vi, vj, vk = T.axis.remap("SSR", [i, j, k])
|
|
with T.init():
|
|
C[vi, vj] = 0.0
|
|
C[vi, vj] = C[vi, vj] + A[vi, vk] * B[vj, vk]
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def element_wise(a: T.handle, c: T.handle) -> None:
|
|
m = T.int32()
|
|
n = T.int32()
|
|
A = T.match_buffer(a, (m, n), "float32")
|
|
C = T.match_buffer(c, (m, n), "float32")
|
|
|
|
B = T.sblock_alloc_buffer((m, n), "float32")
|
|
|
|
for i, j in T.grid(m, n):
|
|
with T.sblock("B"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
B[vi, vj] = A[vi, vj] * 2.0
|
|
|
|
for i, j in T.grid(m, n):
|
|
with T.sblock("C"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
C[vi, vj] = B[vi, vj] + 1.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def element_wise_128_64(a: T.handle, c: T.handle) -> None:
|
|
A = T.match_buffer(a, (128, 64), "float32")
|
|
C = T.match_buffer(c, (128, 64), "float32")
|
|
B = T.sblock_alloc_buffer((128, 64), "float32")
|
|
|
|
for i, j in T.grid(128, 64):
|
|
with T.sblock("B"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
B[vi, vj] = A[vi, vj] * 2.0
|
|
|
|
for i, j in T.grid(128, 64):
|
|
with T.sblock("C"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
C[vi, vj] = B[vi, vj] + 1.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def element_wise_128_n(a: T.handle, c: T.handle) -> None:
|
|
n = T.int32()
|
|
A = T.match_buffer(a, (128, n), "float32")
|
|
C = T.match_buffer(c, (128, n), "float32")
|
|
B = T.sblock_alloc_buffer((128, n), "float32")
|
|
|
|
for i, j in T.grid(128, n):
|
|
with T.sblock("B"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
B[vi, vj] = A[vi, vj] * 2.0
|
|
|
|
for i, j in T.grid(128, n):
|
|
with T.sblock("C"):
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
C[vi, vj] = B[vi, vj] + 1.0
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def mem_copy(a: T.handle, b: T.handle, m: T.int32, n: T.int32, p: T.int32, q: T.int32) -> None:
|
|
A = T.match_buffer(a, (m, n), "float32", strides=[p, 1], elem_offset=q)
|
|
B = T.match_buffer(b, (m, n), "float32", strides=[p, 1], elem_offset=q)
|
|
|
|
for i, j in T.grid(m, n):
|
|
with T.sblock():
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
B[vi, vj] = A[vi, vj]
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def mem_copy_16_16_8_4(a: T.handle, b: T.handle) -> None:
|
|
A = T.match_buffer(a, (16, 16), "float32", strides=[8, 1], elem_offset=4)
|
|
B = T.match_buffer(b, (16, 16), "float32", strides=[8, 1], elem_offset=4)
|
|
|
|
for i, j in T.grid(16, 16):
|
|
with T.sblock():
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
B[vi, vj] = A[vi, vj]
|
|
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def mem_copy_m_n_p_n(a: T.handle, b: T.handle, m: T.int32, n: T.int32, p: T.int32) -> None:
|
|
A = T.match_buffer(a, (m, n), "float32", strides=[p, 1], elem_offset=n)
|
|
B = T.match_buffer(b, (m, n), "float32", strides=[p, 1], elem_offset=n)
|
|
|
|
for i, j in T.grid(m, n):
|
|
with T.sblock():
|
|
vi, vj = T.axis.remap("SS", [i, j])
|
|
B[vi, vj] = A[vi, vj]
|
|
|
|
|
|
def test_specialize_nothing():
|
|
func = matmul.specialize({})
|
|
assert func.same_as(matmul) # Pointer the same
|
|
|
|
|
|
def test_specialize_matmul():
|
|
a, _, _, n = matmul.params
|
|
# fully specialized
|
|
func = matmul.specialize({a: tvm.tirx.decl_buffer((128, 128))})
|
|
assert_structural_equal_ignore_global_symbol(func, matmul_128)
|
|
# partially specialized
|
|
func = matmul.specialize({n: 128})
|
|
assert_structural_equal_ignore_global_symbol(func, matmul_m_128)
|
|
# symbolic specialized
|
|
func = matmul.specialize({n: tvm.tirx.Var("x", "int32") * 8})
|
|
assert_structural_equal_ignore_global_symbol(func, matmul_m_8x)
|
|
|
|
|
|
def test_specialize_elemwise():
|
|
a, c = element_wise.params
|
|
C = element_wise.buffer_map[c]
|
|
# fully specialized
|
|
func = element_wise.specialize({a: tvm.tirx.decl_buffer((128, 64))})
|
|
assert_structural_equal_ignore_global_symbol(func, element_wise_128_64)
|
|
# partially specialized
|
|
func = element_wise.specialize({c: tvm.tirx.decl_buffer((128, C.shape[1]))})
|
|
assert_structural_equal_ignore_global_symbol(func, element_wise_128_n)
|
|
|
|
|
|
def test_specialize_mem_copy():
|
|
a, _, m, n, p, q = mem_copy.params
|
|
# fully specialized
|
|
func = mem_copy.specialize({a: tvm.tirx.decl_buffer((16, 16), strides=[8, 1], elem_offset=4)})
|
|
assert_structural_equal_ignore_global_symbol(func, mem_copy_16_16_8_4)
|
|
func = mem_copy.specialize({n: 16, m: 16, p: 8, q: 4})
|
|
assert_structural_equal_ignore_global_symbol(func, mem_copy_16_16_8_4)
|
|
# partially specialized
|
|
func = mem_copy.specialize({q: n})
|
|
assert_structural_equal_ignore_global_symbol(func, mem_copy_m_n_p_n)
|
|
|
|
|
|
def test_specialize_recursive_load():
|
|
# TODO(Siyuan): add recursive Load testcase, e.g. A[C[i]]
|
|
pass
|
|
|
|
|
|
def test_specialize_with_const_folding():
|
|
@T.prim_func(s_tir=True)
|
|
def before(a: T.handle, b: T.handle):
|
|
n = T.int32()
|
|
A = T.match_buffer(a, [n // 8, 8], "int32")
|
|
B = T.match_buffer(b, [n], "int32")
|
|
for i in range(n - 1):
|
|
with T.sblock():
|
|
vi = T.axis.S(n - 1, i)
|
|
B[vi] = A[vi // 8, vi % 8] + (n + 1) * 42
|
|
|
|
@T.prim_func(s_tir=True)
|
|
def expected(a: T.handle, b: T.handle):
|
|
A = T.match_buffer(a, [2, 8], "int32")
|
|
B = T.match_buffer(b, [16], "int32")
|
|
for i in range(15):
|
|
with T.sblock():
|
|
vi = T.axis.S(15, i)
|
|
B[vi] = A[vi // 8, vi % 8] + 714
|
|
|
|
b = before.params[1]
|
|
after = before.specialize({b: tvm.tirx.decl_buffer([16], dtype="int32")})
|
|
assert_structural_equal_ignore_global_symbol(expected, after)
|
|
|
|
|
|
def test_specialize_decl_buffer():
|
|
"""Buffers occurring in a DeclBuffer statement should be updated"""
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def before(A_data: T.handle("float32"), A_size: T.int32):
|
|
A_buf = T.decl_buffer(A_size, "float32", data=A_data)
|
|
for i in range(A_size):
|
|
A_buf[i] = A_buf[i] * 2.0
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def expected(A_data: T.handle("float32")):
|
|
A_buf = T.decl_buffer(16, "float32", data=A_data)
|
|
for i in range(16):
|
|
A_buf[i] = A_buf[i] * 2.0
|
|
|
|
param_map = {before.params[1]: T.int32(16)}
|
|
after = before.specialize(param_map)
|
|
|
|
tvm.ir.assert_structural_equal(expected, after)
|
|
|
|
|
|
def test_specialize_buffer_var_to_var():
|
|
"""A buffer var may be remapped by specialization
|
|
|
|
If a buffer variable is replaced by a specialization, then other
|
|
buffers using the same buffer var should also be updated.
|
|
"""
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def before(A: T.Buffer([16, 16], "float32"), B: T.Buffer([16, 16], "float32")):
|
|
A_flat = T.decl_buffer([256], "float32", data=A.data)
|
|
B_flat = T.decl_buffer([256], "float32", data=B.data)
|
|
for i in range(256):
|
|
B_flat[i] = A_flat[i] * 2.0
|
|
|
|
# well-formed checker complains about multiple nested definitions of B_flat
|
|
# since it appears in the buffer map twice
|
|
@T.prim_func(private=True, check_well_formed=False, s_tir=True)
|
|
def expected(A: T.Buffer([16, 16], "float32"), B_handle: T.handle):
|
|
B = T.match_buffer(B_handle, [16, 16], "float32", data=A.data)
|
|
A_flat = T.decl_buffer([256], "float32", data=A.data)
|
|
B_flat = T.decl_buffer([256], "float32", data=A.data)
|
|
for i in range(256):
|
|
B_flat[i] = A_flat[i] * 2.0
|
|
|
|
A = before.buffer_map[before.params[0]]
|
|
B_handle = before.params[1]
|
|
param_map = {B_handle: A}
|
|
after = before.specialize(param_map)
|
|
|
|
tvm.ir.assert_structural_equal(expected, after)
|
|
|
|
|
|
def test_specialize_buffer_var_to_expr():
|
|
"""Handle specialization of buffer var
|
|
|
|
The `tirx::Buffer::data` field must be an explicit `tirx::Var`, and
|
|
cannot be replaced with a handle-typed `tirx::Expr`. However,
|
|
these substitutions are useful
|
|
when lowering. If these occur, a binding of the `tirx::Var` is
|
|
included in the specialized function.
|
|
"""
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def before(A_data: T.handle("float32"), B_data: T.handle("float32")):
|
|
A_buf = T.decl_buffer(32, "float32", data=A_data)
|
|
B_buf = T.decl_buffer(16, "float32", data=B_data)
|
|
for i in range(16):
|
|
B_buf[i] = A_buf[i] * 2.0
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def expected(A_data: T.handle("float32")):
|
|
A_buf = T.decl_buffer(32, "float32", data=A_data)
|
|
B_data: T.let[T.Ptr[T.float32]] = T.address_of(A_buf[16])
|
|
B_buf = T.decl_buffer(16, "float32", data=B_data)
|
|
for i in range(16):
|
|
B_buf[i] = A_buf[i] * 2.0
|
|
|
|
B_data = before.params[1]
|
|
# body is a SeqStmt; the first statement is DeclBuffer for A_buf
|
|
A_buf = before.body[0].buffer
|
|
param_map = {B_data: tvm.tirx.address_of(A_buf[16])}
|
|
after = before.specialize(param_map)
|
|
|
|
tvm.ir.assert_structural_equal(expected, after)
|
|
|
|
|
|
def test_specialization_updates_ty():
|
|
"""Update type in specialization
|
|
|
|
A PrimFunc may have a `relax.Type`. If that PrimFunc is
|
|
specialized, the type should be updated.
|
|
"""
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def before(n: T.int32) -> T.int32:
|
|
T.ret(n * 10)
|
|
|
|
@T.prim_func(private=True, s_tir=True)
|
|
def expected() -> T.int32:
|
|
T.ret(50)
|
|
|
|
ty_before = tvm.relax.FuncType([tvm.ir.PrimType("int32")], tvm.ir.PrimType("int32"))
|
|
tvm.ir.assert_structural_equal(before.ty, ty_before)
|
|
|
|
ty_expected = tvm.relax.FuncType([], tvm.ir.PrimType("int32"))
|
|
tvm.ir.assert_structural_equal(expected.ty, ty_expected)
|
|
|
|
n = before.params[0]
|
|
param_map = {n: 5}
|
|
after = before.specialize(param_map)
|
|
|
|
tvm.ir.assert_structural_equal(after, expected)
|
|
tvm.ir.assert_structural_equal(after.ty, ty_expected)
|
|
|
|
|
|
if __name__ == "__main__":
|
|
tvm.testing.main()
|