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 9999114e70 [Codegen][OpenCL] fix amibiguous selection operator call
(#14833)
9999114e70 is described below
commit 9999114e70c5587edd2bf0b14c9b384f843b849d
Author: Bohan Hou <[email protected]>
AuthorDate: Sun May 14 18:42:07 2023 -0700
[Codegen][OpenCL] fix amibiguous selection operator call (#14833)
* fix
* fix
* Update test_target_codegen_opencl.py
* Update test_target_codegen_opencl.py
* Update test_target_codegen_opencl.py
* Update codegen_opencl.cc
---------
Co-authored-by: Tianqi Chen <[email protected]>
---
src/target/source/codegen_opencl.cc | 6 +---
.../python/unittest/test_target_codegen_opencl.py | 32 ++++++++++++++++------
2 files changed, 24 insertions(+), 14 deletions(-)
diff --git a/src/target/source/codegen_opencl.cc
b/src/target/source/codegen_opencl.cc
index de96f923e2..14fbc4d840 100644
--- a/src/target/source/codegen_opencl.cc
+++ b/src/target/source/codegen_opencl.cc
@@ -567,11 +567,7 @@ void CodeGenOpenCL::VisitExpr_(const SelectNode* op,
std::ostream& os) {
os << ", ";
PrintExpr(op->condition, oss);
if (op->dtype.is_float()) {
- if (op->condition.dtype().is_uint() || op->condition.dtype().is_int()) {
- os << oss.str();
- } else {
- os << CastTo(oss.str(), DataType::Int(op->dtype.bits(),
op->dtype.lanes()));
- }
+ os << CastTo(oss.str(), DataType::Int(op->dtype.bits(),
op->dtype.lanes()));
} else {
os << CastFromTo(oss.str(), op->condition.dtype(), op->dtype);
}
diff --git a/tests/python/unittest/test_target_codegen_opencl.py
b/tests/python/unittest/test_target_codegen_opencl.py
index bc2d0a84fd..67dc37363e 100644
--- a/tests/python/unittest/test_target_codegen_opencl.py
+++ b/tests/python/unittest/test_target_codegen_opencl.py
@@ -168,20 +168,34 @@ def test_opencl_type_casting():
c = tvm.nd.empty((n,), dtype, ctx)
assembly = fun.imported_modules[0].get_source()
- false_branch = "((float4)(0.000000e+00f, 0.000000e+00f, 0.000000e+00f,
0.000000e+00f))"
- true_branch = "((float4)(1.000000e+00f, 1.000000e+00f, 1.000000e+00f,
1.000000e+00f))"
- lcond = "(convert_uint4(((uint4)((((int)get_local_id(0)) == 3),
(((int)get_local_id(0)) == 3), (((int)get_local_id(0)) == 3),
(((int)get_local_id(0)) == 3)))))"
- rcond = "(convert_uint4((((int4)((0)+(1*0), (0)+(1*1), (0)+(1*2),
(0)+(1*3))) == ((int4)(3, 3, 3, 3)))))"
- cond = "({} && {})".format(lcond, rcond)
- select = "select({}, {}, {})".format(false_branch, true_branch, cond)
- count = assembly.count(select)
- assert count == 1
- fun(c)
+ if dtype == "float32":
+ false_branch = "((float4)(0.000000e+00f, 0.000000e+00f,
0.000000e+00f, 0.000000e+00f))"
+ true_branch = "((float4)(1.000000e+00f, 1.000000e+00f,
1.000000e+00f, 1.000000e+00f))"
+ lcond =
"convert_int4(((convert_uint4(((uint4)((((int)get_local_id(0)) == 3),
(((int)get_local_id(0)) == 3), (((int)get_local_id(0)) == 3),
(((int)get_local_id(0)) == 3)))))"
+ rcond = "(convert_uint4((((int4)((0)+(1*0), (0)+(1*1), (0)+(1*2),
(0)+(1*3))) == ((int4)(3, 3, 3, 3)))))"
+ cond = "({} && {})".format(lcond, rcond)
+ select = "select({}, {}, {})".format(false_branch, true_branch,
cond)
+ count = assembly.count(select)
+ assert count == 1
+ fun(c)
+
+ elif dtype == "float16":
+ false_branch = "((half4)((half)0.000000e+00f, (half)0.000000e+00f,
(half)0.000000e+00f, (half)0.000000e+00f))"
+ true_branch = "((half4)((half)1.000000e+00f, (half)1.000000e+00f,
(half)1.000000e+00f, (half)1.000000e+00f))"
+ lcond =
"convert_short4(((convert_uint4(((uint4)((((int)get_local_id(0)) == 3),
(((int)get_local_id(0)) == 3), (((int)get_local_id(0)) == 3),
(((int)get_local_id(0)) == 3)))))"
+ rcond = "(convert_uint4((((int4)((0)+(1*0), (0)+(1*1), (0)+(1*2),
(0)+(1*3))) == ((int4)(3, 3, 3, 3)))))))"
+ cond = "({} && {})".format(lcond, rcond)
+ select = "select({}, {}, {})".format(false_branch, true_branch,
cond)
+ count = assembly.count(select)
+ assert count == 1
+ fun(c)
dev = tvm.device(target, 0)
check_type_casting(dev, 16, "float32")
+ # fp16 is not yet supported in ci
+ # check_type_casting(dev, 16, "float16")
if __name__ == "__main__":
