From dadbbec0f59738a425d8982cecb257ce83580c78 Mon Sep 17 00:00:00 2001
From: Egor Churaev <egor.churaev@gmail.com>
Date: Thu, 8 Jun 2023 11:44:27 +0300
Subject: [PATCH] [TOPI][Adreno] Fix problem with ceil_log2

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 +++
 .../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 4d305d1e2f2e..8b66ca2cc9ea 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 4a426c952b56..83612b7f5979 100644
--- a/tests/python/unittest/test_target_codegen_opencl.py
+++ b/tests/python/unittest/test_target_codegen_opencl.py
@@ -190,5 +190,32 @@ def check_type_casting(ctx, n, dtype):
     # check_type_casting(dev, 16, "float16")
 
 
+@tvm.testing.requires_gpu
+@tvm.testing.requires_opencl
+@tvm.testing.parametrize_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()