Check failed: ret == 0 (-1 vs. 0) : TVMError: OpenCL build error for device=0x79f8964678 BC-src-code:4:521: error: expected ')'

Questions
#1

my tvm version is 0.8.dev0

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
Relay opencl quantization error
#2

my deploy cpp code：

cc_code
cc_code928×927 52 KB

compile command

compile
compile865×163 12 KB

#3

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))))))))));
    }
  }
}
#4

Hi, have you find the solution？the same error occurs when i use tvm auto schedule to tune a operator on mobile gpu(opencl).