This is an automated email from the ASF dual-hosted git repository.
masahi 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 df5b180ebd [TOPI][Adreno] Fix problem with ceil_log2 (#15061)
df5b180ebd is described below
commit df5b180ebd31dd2be5e289adae52707a5bb6dd21
Author: Egor Churaev <[email protected]>
AuthorDate: Fri Jun 9 19:28:50 2023 +0300
[TOPI][Adreno] Fix problem with ceil_log2 (#15061)
On Adreno devices double type is not supported by OpenCL. Modified
`ceil_log2` operation to cast value to float instead of double in case
of Adreno.
---
python/tvm/topi/math.py | 3 +++
.../python/unittest/test_target_codegen_opencl.py | 27 ++++++++++++++++++++++
2 files changed, 30 insertions(+)
diff --git a/python/tvm/topi/math.py b/python/tvm/topi/math.py
index 4d305d1e2f..8b66ca2cc9 100644
--- a/python/tvm/topi/math.py
+++ b/python/tvm/topi/math.py
@@ -865,4 +865,7 @@ def ceil_log2(x):
return res
+ if "adreno" in tvm.target.Target.current().device_name:
+ return cast(tvm.tir.ceil(tvm.tir.log2(cast(x, "float32"))), x.dtype)
+
return cast(tvm.tir.ceil(tvm.tir.log2(cast(x, "float64"))), x.dtype)
diff --git a/tests/python/unittest/test_target_codegen_opencl.py
b/tests/python/unittest/test_target_codegen_opencl.py
index 4a426c952b..83612b7f59 100644
--- a/tests/python/unittest/test_target_codegen_opencl.py
+++ b/tests/python/unittest/test_target_codegen_opencl.py
@@ -190,5 +190,32 @@ def test_opencl_type_casting():
# check_type_casting(dev, 16, "float16")
[email protected]_gpu
[email protected]_opencl
[email protected]_targets("opencl", "opencl -device=adreno")
+def test_opencl_ceil_log2(target):
+ def _check(target, n, dtype):
+ with tvm.target.Target(target):
+ C = te.compute(
+ (n,),
+ lambda i: tvm.topi.ceil_log2(i),
+ name="C",
+ )
+ func = te.create_prim_func([C])
+ sch = tvm.tir.Schedule(func)
+ (x,) = sch.get_loops(sch.get_block("C"))
+ sch.bind(x, "threadIdx.x")
+
+ fun = tvm.build(sch.mod, target=target)
+ assembly = fun.imported_modules[0].get_source()
+ if "adreno" in target:
+ pattern = "convert_float"
+ else:
+ pattern = "convert_double"
+ assert assembly.count(pattern) != 0
+
+ _check(target, 32, "float32")
+
+
if __name__ == "__main__":
tvm.testing.main()
