I need to implement opencl operator resize, I modified the function
check_target() in tests/python/topi/python/test_topi_image.py,The program can
be executed normally.
arch = "arm64"
target_host = "llvm -mtriple=%s-linux-android" % arch
target = tvm.target.Target("opencl", host="llvm
-mtriple=arm64-linux-android")
def check_target(target):
print("Running on target: %s" % target)
with tvm.target.Target(target):
s = tvm.topi.testing.get_injective_schedule(target)(B)
# Establish remote connection with target hardware
f = tvm.build(s, [A, B], target, name="resize2d")
#temp = utils.tempdir()
#path_dso_cl = temp.relpath("dev_lib_cl.so")
path_dso_cl = '../howto_deploy/dev_lib_cl.so'
print(path_dso_cl)
f.export_library(path_dso_cl, ndk.create_shared)
print("init RPC ")
tracker = rpc.connect_tracker(tracker_host, tracker_port)
remote = tracker.request(key, priority=0, session_timeout=320)
remote.upload(path_dso_cl)
dev = remote.cl()
print("Run GPU(OpenCL Flavor) test ...")
print(dev)
f = remote.load_module("dev_lib_cl.so")
a = tvm.nd.array(a_np, dev)
b = tvm.nd.array(np.zeros(out_shape, dtype=dtype), dev)
f(a, b)
#tvm.testing.assert_allclose(b.numpy(), b_np, rtol=1e-3, atol=1e-3)
#time cast
time_f = f.time_evaluator(f.entry_name, dev, number=10)
cost = time_f(a, b).mean
print("%g secs/op\n" % cost)
evaluator = f.time_evaluator(f.entry_name, dev, number=10,
min_repeat_ms=500)
print(
"Execution time of this operator: %.3f ms"
% (np.median(evaluator(a, b).results) * 1000)
)
I modified apps/howto_deploy/cpp_deploy.cc
The code in and use the following method to compile, crsh appears when executed
on the phone.
` terminating with uncaught exception of type
tvm::runtime::InternalError: [14:17:41]
/workspace/source/tvm/src/./runtime/library_module.cc:78:`
---------------------------------------------------------------
An error occurred during the execution of TVM.
For more information, please see:
https://tvm.apache.org/docs/errors.html
---------------------------------------------------------------
Check failed: ret == 0 (-1 vs. 0) : TVMError: OpenCL build error for
device=0x737a0b8678
BC-src-code:4:521: error: expected ')'
resize[((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))))] =
((uchar)((((((f_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) / 6048))))), 519), 0)
* 756) + (max(mied_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) +
((int)get_local_id(0))) % 6048) / 3))))), 251), 0) <U+0001>d<C8>)) +
((((i0_i10)) * 256)) + ((int)get_local_id(0))) % 3)))]) * (1.000000e+00f -
((1.250000e-01f * ((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536)
6048) / 3))) - ((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 25A[((((max(min(((int)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((intint)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0_fused_i3_fused_outer *
65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) *
((1.250000e-01f * ((float)(((((i0_i1_fu * 256)) + ((int)get_local_id(0))) %
6048) / 3))) - ((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 6))) % 6048) / 3))))))))) *
(1.000000e+00f - ((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(t)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0)))0000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) / 6048)))))
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 6048) / 3))))),
251), (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) *
(1.000000e+00f - ((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3t)get_local_id(0))) % 6048) / 3))) -
((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)))))))))) +
(((float)A[((((max(min((((int)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_),
0) * 756) + (max(min((((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 2 * 3)) + ((((i0_i1_fused_i2_fused_i3_fused_outer *
65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) *
((1.250000e-) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) %
6048) / 3))) - ((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_56)) + ((int)get_local_id(0))) % 6048) / 3))))))))) *
((1.250000e-01f * ((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)getfloat)((int)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_loc
^
BC-src-code:4:338: note: to match this '('
resize[((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))))] =
((uchar)((((((f_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) / 6048))))), 519), 0)
* 756) + (max(mied_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) +
((int)get_local_id(0))) % 6048) / 3))))), 251), 0) <U+0001>d<C8>)) +
((((i0_i10)) * 256)) + ((int)get_local_id(0))) % 3)))]) * (1.000000e+00f -
((1.250000e-01f * ((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536)
6048) / 3))) - ((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 25A[((((max(min(((int)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((intint)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0_fused_i3_fused_outer *
65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) *
((1.250000e-01f * ((float)(((((i0_i1_fu * 256)) + ((int)get_local_id(0))) %
6048) / 3))) - ((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 6))) % 6048) / 3))))))))) *
(1.000000e+00f - ((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(t)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0)))0000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) / 6048)))))
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 6048) / 3))))),
251), (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) *
(1.000000e+00f - ((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3t)get_local_id(0))) % 6048) / 3))) -
((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)))))))))) +
(((float)A[((((max(min((((int)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_),
0) * 756) + (max(min((((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 2 * 3)) + ((((i0_i1_fused_i2_fused_i3_fused_outer *
65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) *
((1.250000e-) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) %
6048) / 3))) - ((float)((int)floor((1.250000e-01f *
((float)(((((i0_i1_fused_56)) + ((int)get_local_id(0))) % 6048) / 3))))))))) *
((1.250000e-01f * ((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)getfloat)((int)floor((1.250000e-01f *
((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) +
(((int)get_group_id(0)) * 256)) + ((int)get_loc
^
1 diagnostic(s) generated.
Stack trace:
File "/workspace/source/tvm/src/./runtime/opencl/opencl_module.cc", line
244
Stack trace not available when DMLC_LOG_STACK_TRACE is disabled at compile
time.
Aborted
---
[Visit
Topic](https://discuss.tvm.apache.org/t/check-failed-ret-0-1-vs-0-tvmerror-opencl-build-error-for-device-0x79f8964678-bc-src-code521-error-expected/10838/1)
to respond.
You are receiving this because you enabled mailing list mode.
To unsubscribe from these emails, [click
here](https://discuss.tvm.apache.org/email/unsubscribe/33f24001b3b40870d94357dc314426fd10c3e9a3d97ef77fb0132ae968f6f787).
