You are viewing a plain text version of this content. The canonical link for it is here.
Posted to discuss-archive@tvm.apache.org by donglaxiche via Apache TVM Discuss <no...@discuss.tvm.ai> on 2021/08/19 07:42:48 UTC
[Apache TVM Discuss] [Questions] Check failed: ret == 0 (-1 vs. 0) :
TVMError: OpenCL build error for device=0x79f8964678 BC-src-code:4:521: error:
expected ')'
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).
[Apache TVM Discuss] [Questions] Check failed: ret == 0 (-1 vs. 0) :
TVMError: OpenCL build error for device=0x79f8964678 BC-src-code:4:521: error:
expected ')'
Posted by donglaxiche via Apache TVM Discuss <no...@discuss.tvm.ai>.
my deploy cpp code：
compile command
---
[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/2) 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/940163a7a905d4936aae1720de2f1a649873cb844f3cad9c86bd95f3610f7941).
[Apache TVM Discuss] [Questions] Check failed: ret == 0 (-1 vs. 0) : TVMError: OpenCL build error for device=0x79f8964678 BC-src-code:4:521: error: expected ')'
Posted by Zhangwenzheng via Apache TVM Discuss <no...@discuss.tvm.ai>.
Hi, have you find the solution？the same error occurs when i use tvm auto schedule to tune a operator on mobile gpu(opencl).
---
[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/4) 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/65f7211e0df4af4f1bbba636da3b243c30301d5fdcb14712cf442c06cae17c30).
[Apache TVM Discuss] [Questions] Check failed: ret == 0 (-1 vs. 0) :
TVMError: OpenCL build error for device=0x79f8964678 BC-src-code:4:521: error:
expected ')'
Posted by donglaxiche via Apache TVM Discuss <no...@discuss.tvm.ai>.
i have print kennel in opencl_module.cc，It looks like there is a problem of symbol mismatch。
__kernel void resize2d_kernel0(__global uchar* restrict resize, __global uchar* restrict A) {
for (int i0_i1_fused_i2_fused_i3_fused_outer = 0; i0_i1_fused_i2_fused_i3_fused_outer < 384; ++i0_i1_fused_i2_fused_i3_fused_outer) {
if ((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) < 25159680) {
resize[((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))))] = ((uchar)((((((float)A[((((max(min(((int)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))) / 6048))))), 519), 0) * 756) + (max(min(((int)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))) % 6048) / 3))))), 251), 0) dƩ) + ((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) * (1.000000e+00f - ((1.250000e-01f * ((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 6048) / 3))) - ((float)((int)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))) % 6048) / 3))))))))) + (((float)A[((((max(min(((int)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))) / 6048))))), 519), 0) * 756) + (max(min((((int)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))) % 6048) / 3))))) + 1), 251), 0) * 3)) + ((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) * ((1.250000e-01f * ((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 6048) / 3))) - ((float)((int)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))) % 6048) / 3))))))))) * (1.000000e+00f - ((1.250000e-01f * ((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) / 6048))) - ((float)((int)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))) / 6048))))))))) + (((((float)A[((((max(min((((int)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))) / 6048))))) + 1), 519), 0) * 756) + (max(min(((int)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))) % 6048) / 3))))), 251), 0) * 3)) + ((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) * (1.000000e+00f - ((1.250000e-01f * ((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 6048) / 3))) - ((float)((int)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))) % 6048) / 3))))))))) + (((float)A[((((max(min((((int)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))) / 6048))))) + 1), 519), 0) * 756) + (max(min((((int)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))) % 6048) / 3))))) + 1), 251), 0) * 3)) + ((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 3)))]) * ((1.250000e-01f * ((float)(((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) % 6048) / 3))) - ((float)((int)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))) % 6048) / 3))))))))) * ((1.250000e-01f * ((float)((((i0_i1_fused_i2_fused_i3_fused_outer * 65536) + (((int)get_group_id(0)) * 256)) + ((int)get_local_id(0))) / 6048))) - ((float)((int)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))) / 6048))))))))));
}
}
}
---
[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/3) 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/e642a38a0e06551a7d9996e6860aa92990ea267facf53c786a92ae3ffa5cbe96).