You are viewing a plain text version of this content. The canonical link for it is here.
Posted to commits@tvm.apache.org by tq...@apache.org on 2021/02/02 17:51:46 UTC
[tvm] branch main updated: Fix missing round(), floor(),
ceil() for target C lowering (#7382)
This is an automated email from the ASF dual-hosted git repository.
tqchen 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 de0ab4c Fix missing round(), floor(), ceil() for target C lowering (#7382)
de0ab4c is described below
commit de0ab4cb7cd73e3c05eb3512c916607587156b13
Author: Balint Cristian <cr...@gmail.com>
AuthorDate: Tue Feb 2 19:51:31 2021 +0200
Fix missing round(), floor(), ceil() for target C lowering (#7382)
---
src/target/intrin_rule.cc | 6 ++
.../python/unittest/test_target_codegen_c_host.py | 87 ++++++++++++++++++++--
2 files changed, 87 insertions(+), 6 deletions(-)
diff --git a/src/target/intrin_rule.cc b/src/target/intrin_rule.cc
index f8f4d0e..1a72144 100644
--- a/src/target/intrin_rule.cc
+++ b/src/target/intrin_rule.cc
@@ -77,6 +77,12 @@ TVM_REGISTER_GLOBAL("tvm.intrin.rule.default.ldexp").set_body(DispatchPureExtern
TVM_REGISTER_GLOBAL("tvm.intrin.rule.default.sqrt").set_body(DispatchPureExtern<FloatSuffix>);
+TVM_REGISTER_GLOBAL("tvm.intrin.rule.default.floor").set_body(DispatchPureExtern<FloatSuffix>);
+
+TVM_REGISTER_GLOBAL("tvm.intrin.rule.default.ceil").set_body(DispatchPureExtern<FloatSuffix>);
+
+TVM_REGISTER_GLOBAL("tvm.intrin.rule.default.round").set_body(DispatchPureExtern<FloatSuffix>);
+
TVM_REGISTER_GLOBAL("tvm.intrin.rule.default.rsqrt")
.set_body([](const TVMArgs& args, TVMRetValue* rv) {
PrimExpr e = args[0];
diff --git a/tests/python/unittest/test_target_codegen_c_host.py b/tests/python/unittest/test_target_codegen_c_host.py
index 3178d6d..d1ca8b1 100644
--- a/tests/python/unittest/test_target_codegen_c_host.py
+++ b/tests/python/unittest/test_target_codegen_c_host.py
@@ -30,12 +30,12 @@ def test_add():
s = te.create_schedule(C.op)
def check_c():
- mhost = tvm.build(s, [A, B, C], "c", name="fadd")
+ mhost = tvm.build(s, [A, B, C], "c", name="test_fadd")
temp = utils.tempdir()
path_dso = temp.relpath("temp.so")
mhost.export_library(path_dso)
m = tvm.runtime.load_module(path_dso)
- fadd = m["fadd"]
+ fadd = m["test_fadd"]
ctx = tvm.cpu(0)
# launch the kernel.
n = nn
@@ -73,14 +73,14 @@ def test_add_pipeline():
)
binds = {A: Ab}
# BUILD and invoke the kernel.
- f1 = tvm.lower(s, [A, B, C], name="fadd_pipeline")
+ f1 = tvm.lower(s, [A, B, C], name="test_fadd_pipeline")
mhost = tvm.build(f1, target="c")
temp = utils.tempdir()
path_dso = temp.relpath("temp.so")
mhost.export_library(path_dso)
m = tvm.runtime.load_module(path_dso)
- fadd = m["fadd_pipeline"]
+ fadd = m["test_fadd_pipeline"]
ctx = tvm.cpu(0)
# launch the kernel.
n = nn
@@ -103,12 +103,12 @@ def test_reinterpret():
s = te.create_schedule(B.op)
def check_c():
- mhost = tvm.build(s, [A, B], "c", name="reinterpret")
+ mhost = tvm.build(s, [A, B], "c", name="test_reinterpret")
temp = utils.tempdir()
path_dso = temp.relpath("temp.so")
mhost.export_library(path_dso)
m = tvm.runtime.load_module(path_dso)
- fadd = m["reinterpret"]
+ fadd = m["test_reinterpret"]
ctx = tvm.cpu(0)
n = nn
a = tvm.nd.array(np.random.randint(-(2 ** 30), 2 ** 30, size=n).astype(A.dtype), ctx)
@@ -119,7 +119,82 @@ def test_reinterpret():
check_c()
+def test_ceil():
+ nn = 1024
+ n = tvm.runtime.convert(nn)
+ A = te.placeholder((n,), name="A", dtype="float32")
+ B = te.compute(A.shape, lambda *i: tvm.tir.call_intrin("float32", "tir.ceil", A(*i)), name="B")
+ s = te.create_schedule(B.op)
+
+ def check_c():
+ mhost = tvm.build(s, [A, B], "c", name="test_ceil")
+ temp = utils.tempdir()
+ path_dso = temp.relpath("temp.so")
+ mhost.export_library(path_dso)
+ m = tvm.runtime.load_module(path_dso)
+ fceil = m["test_ceil"]
+ ctx = tvm.cpu(0)
+ n = nn
+ a = tvm.nd.array(np.random.rand(n).astype(A.dtype), ctx)
+ b = tvm.nd.array(np.zeros(n, dtype=B.dtype), ctx)
+ fceil(a, b)
+ tvm.testing.assert_allclose(b.asnumpy(), (np.ceil(a.asnumpy()).view("float32")))
+
+ check_c()
+
+
+def test_floor():
+ nn = 1024
+ n = tvm.runtime.convert(nn)
+ A = te.placeholder((n,), name="A", dtype="float32")
+ B = te.compute(A.shape, lambda *i: tvm.tir.call_intrin("float32", "tir.floor", A(*i)), name="B")
+ s = te.create_schedule(B.op)
+
+ def check_c():
+ mhost = tvm.build(s, [A, B], "c", name="test_floor")
+ temp = utils.tempdir()
+ path_dso = temp.relpath("temp.so")
+ mhost.export_library(path_dso)
+ m = tvm.runtime.load_module(path_dso)
+ ffloor = m["test_floor"]
+ ctx = tvm.cpu(0)
+ n = nn
+ a = tvm.nd.array(np.random.rand(n).astype(A.dtype), ctx)
+ b = tvm.nd.array(np.zeros(n, dtype=B.dtype), ctx)
+ ffloor(a, b)
+ tvm.testing.assert_allclose(b.asnumpy(), (np.floor(a.asnumpy()).view("float32")))
+
+ check_c()
+
+
+def test_round():
+ nn = 1024
+ n = tvm.runtime.convert(nn)
+ A = te.placeholder((n,), name="A", dtype="float32")
+ B = te.compute(A.shape, lambda *i: tvm.tir.call_intrin("float32", "tir.round", A(*i)), name="B")
+ s = te.create_schedule(B.op)
+
+ def check_c():
+ mhost = tvm.build(s, [A, B], "c", name="test_round")
+ temp = utils.tempdir()
+ path_dso = temp.relpath("temp.so")
+ mhost.export_library(path_dso)
+ m = tvm.runtime.load_module(path_dso)
+ fround = m["test_round"]
+ ctx = tvm.cpu(0)
+ n = nn
+ a = tvm.nd.array(np.random.rand(n).astype(A.dtype), ctx)
+ b = tvm.nd.array(np.zeros(n, dtype=B.dtype), ctx)
+ fround(a, b)
+ tvm.testing.assert_allclose(b.asnumpy(), (np.round(a.asnumpy()).view("float32")))
+
+ check_c()
+
+
if __name__ == "__main__":
test_add()
test_add_pipeline()
test_reinterpret()
+ test_ceil()
+ test_floor()
+ test_round()