openclc to llvm

Hi all,

I am trying to generate an llvm file from opencl kernel. It recognizes address space qualifiers for openCL like __global. However I get following errors for the types ( int4 etc… for OpenCL)

– My Simple Kernel —

__kernel void VectorAdd(__global int4 *c, __global int4 *a, __global int4 *b){
unsigned int idx = get_global_id(0);
c[idx] = a[idx] + b[idx];
}

– Compilation and errors

…/build/Debug+Asserts/bin/clang -x cl -emit-llvm -S vectoradd_post.cl -o vectoradd_post.ll
vectoradd_post.cl:1:34: error: unknown type name ‘int4’; did you mean ‘int’?
__kernel void VectorAdd(__global int4 *c, __global int4 *a, __global int4 *b){
^
vectoradd_post.cl:1:52: error: unknown type name ‘int4’; did you mean ‘int’?
__kernel void VectorAdd(__global int4 *c, __global int4 *a, __global int4 *b){
^
vectoradd_post.cl:1:70: error: unknown type name ‘int4’; did you mean ‘int’?
__kernel void VectorAdd(__global int4 *c, __global int4 *a, __global int4 *b){
^
vectoradd_post.cl:2:24: warning: implicit declaration of function ‘get_global_id’ is invalid in C99 [-Wimplicit-function-declaration]
unsigned int idx = get_global_id(0);
^
1 warning and 3 errors generated.

.

Please let me what am I doing wrong. Also is there any document regarding opencl with clang.

Thanks and Regards,
Ankur

Hi Ankur,

“int4” is not a valid C/C++ type on its own. Are you missing a #include that typedefs those types, perhaps?

-Jim

Hi Jim,

Thanks for reply. int4 is a valid data type of Opencl-C language. I read
that OpenCL support is available in Clang with "-x cl" option.

I also found from some conversation on the llvm mailing group that libclc
is also required for supporting builtin functions for openCL specification.
Is it necessary for parsing the OpenCL introduced data types like int4 too?

Please help me with the issue.

Regards,
Ankur

Hi Ankur,

Yes, you need the support of libclc in order to successfully compile OpenCL
kernels. It is available at http://libclc.llvm.org/. After building libclc,
you can use the script "compile-test.sh" to compile OpenCL kernels. For me,
the compile-test.sh has a little problem (due to the option -Xclang
ptx32--nvidiacl/lib/builtins.bc), so I just removed the option, and the
resulting script looks like this:

clang -ccc-host-triple ptx32--nvidiacl -Iptx-nvidiacl/include
-Igeneric/include -include clc/clc.h -target-feature -Xclang +ptx23 -Xclang
-target-feature -Xclang +sm20 -Dcl_clang_storage_class_specifiers -S
-emit-llvm "$@"

Using the modified script, OpenCL kernels can be successfully compiled into
PTX codes.

Have a look at this post:
http://lists.cs.uiuc.edu/pipermail/llvmdev/2011-December/045983.html, it
maybe helpful as well.

I'm also quite new to Clang/LLVM, hope it will solve your problem.

Regards,
Lei Mou

Hi Ankur,

Yes, you need the support of libclc in order to successfully compile OpenCL
kernels. It is available at http://libclc.llvm.org/. After building libclc,
you can use the script "compile-test.sh" to compile OpenCL kernels. For me,
the compile-test.sh has a little problem (due to the option -Xclang
ptx32--nvidiacl/lib/builtins.bc), so I just removed the option, and the

What problem? Did you follow the instructions in README.TXT to build
the builtins.bc file?

resulting script looks like this:

clang -ccc-host-triple ptx32--nvidiacl -Iptx-nvidiacl/include
-Igeneric/include -include clc/clc.h -target-feature -Xclang +ptx23 -Xclang
-target-feature -Xclang +sm20 -Dcl_clang_storage_class_specifiers -S
-emit-llvm "$@"

Using the modified script, OpenCL kernels can be successfully compiled into
PTX codes.

But most of the builtin functions won't work.

Thanks,

Hi Peter,

> Hi Ankur,
>
> Yes, you need the support of libclc in order to successfully compile
OpenCL
> kernels. It is available at http://libclc.llvm.org/. After building
libclc,
> you can use the script "compile-test.sh" to compile OpenCL kernels. For
me,
> the compile-test.sh has a little problem (due to the option -Xclang
> ptx32--nvidiacl/lib/builtins.bc), so I just removed the option, and the

What problem? Did you follow the instructions in README.TXT to build
the builtins.bc file?

Thank you for your attention. I did follow the instructions in README.txt.
I got a lot of errors and warnings like the following one:

ptx32--nvidiacl/lib/builtins.bc:299:196: warning: null character ignored
2�>���M 0 !����<� �D��#>R��� hS � �>b� ODH>�#��;<F PT�5ن �T��4��#��!
!��p@�H}�e�Q ��#� �D��#>���c� HE\��� $ !���-$ �,>�#� �,>�# �T��4�d���
�8
���A8��T>��#un� �

                                           ^
fatal error: too many errors emitted, stopping now [-ferror-limit=]
2900 warnings and 20 errors generated.

I don't know what's the problem, so I just removed the option from the
script.

> resulting script looks like this:
>
> clang -ccc-host-triple ptx32--nvidiacl -Iptx-nvidiacl/include
> -Igeneric/include -include clc/clc.h -target-feature -Xclang +ptx23
-Xclang
> -target-feature -Xclang +sm20 -Dcl_clang_storage_class_specifiers -S
> -emit-llvm "$@"
>
> Using the modified script, OpenCL kernels can be successfully compiled
into
> PTX codes.

But most of the builtin functions won't work.

I only tried a few kernels, including VectorAdd and Matrix Vector
Multiplication in the Nvidia OpenCL SDK. For the kernel mentioned in
Ankur's mail, it works as expected.
Please let me know how can the problem be solved. Thank you very much!

Thanks,
--
Peter

Regards,
Lei Mou