Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[BUG][CODEGEN][OPENCL]Invalid data type in generated OpenCL kernel #5710

Closed
CloudManX opened this issue Jun 1, 2020 · 4 comments
Closed

Comments

@CloudManX
Copy link

Hello all,

I believe there is a bug within OpenCL codegen which results in a typo in data type identifiers. I am recently observing undefined type identifiers such as 'float44' and 'float88' in generated OpenCL codes. Here is an example.


vstore4(((float44)(((__global float*)placeholder)[_1.s0],((__global float*)placeholder)[_1.s1],((__global float*)placeholder)[_1.s2],((__global float*)placeholder)[_1.s3])), 0, (__global float*)kernel_vec + ((((int)get_group_id(0)) * 1024) + (((int)get_local_id(0)) * 4)));

I also found the similar problem in this post https://discuss.tvm.ai/t/bug-tvm-generate-wrong-opencl-code/6658. Here is my code to replicate this error:


import numpy as np
import topi
import tvm 

target = tvm.target.create("opencl -device=mali")
target_host = 'llvm'

dshape = (1, 64, 224, 224)
kshape = (64, 64, 3, 3)

data = tvm.te.placeholder(dshape)
kernel = tvm.te.placeholder(kshape)

with target:
    conv = topi.mali.conv2d_nchw_spatial_pack(data, kernel, [1,1], [1,1,1,1], [1,1],"float32" )
    sch = topi.mali.schedule_conv2d_nchw_spatial_pack([conv])
    mod = tvm.build(sch, [data, kernel], target, target_host )
    print("------opencl code------")
    print(mod.imported_modules[0].get_source())

The printed code will contain a 'float44' as shown in the example above.

In addtion, my investigation of the cause leads to commit d2e58ad
https://github.com/apache/incubator-tvm/blob/d2e58ad2fd92c09dffe064e4efbfc484cf2de6e5/src/target/source/codegen_c.cc#L963

@tqchen @huochaitiantang @kevinthesun Could you please help to confirm if "t.lanes()" is unnecessary here.

Thank you!

@tqchen
Copy link
Member

tqchen commented Jun 1, 2020

cc @kazum @kevinthesun @huochaitiantang@vinx13 please see if you can send a fix.

@vinx13
Copy link
Member

vinx13 commented Jun 3, 2020

I think printing t.lanes() here is unnecessary, a patch is welcomed

@abergeron
Copy link
Contributor

This seems similar to what I've seen in #5704.

@tqchen
Copy link
Member

tqchen commented Jun 3, 2020

fixed by #5722 thanks to @abergeron

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
None yet
Projects
None yet
Development

No branches or pull requests

4 participants