[cfe-dev] builtins name mangling in SPIR 2.0

[Hongbin Zheng via cfe-dev]

Thanks a lot.

On Mon, Sep 12, 2016 at 1:42 PM, Liu, Yaxun (Sam) <Yaxun.Liu at amd.com> wrote:

> If you use the default header file under clang/lib/Headers/opencl-c.h,
> get_global_id will be mangled.
>
>
>
> If you want to declare get_global_id in your own header, add
> __attribute__((overloadable)), then it will be mangled.
>
>
>
> Sam
>
>
>
> *From:* Hongbin Zheng [mailto:etherzhhb at gmail.com]
> *Sent:* Monday, September 12, 2016 4:21 PM
> *To:* cfe-dev at lists.llvm.org; llvm-dev <llvm-dev at lists.llvm.org>; Liu,
> Yaxun (Sam) <Yaxun.Liu at amd.com>
> *Subject:* builtins name mangling in SPIR 2.0
>
>
>
> Hi all,
>
>
>
> According to the SPIR 2.0 spec[1], the name of OpenCL builtins are mangled.
>
>
>
> However, when I compile OpenCl code with Clang 3.9 with the
> "spir64-unknown-unknown" target, Clang generates IR without mangling the
> builtins, e.g. for:
>
>
>
> __kernel void input_zip_int(__global int *in0) {
>
>   *in0 = get_global_id(0);
>
> }
>
>
>
> clang generates:
>
>
>
> define spir_kernel void @input_zip_int(i32 addrspace(1)* nocapture %in0)
> local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4
> !kernel_arg_type !5 !kernel_arg_base_type !5 !kernel_arg_type_qual !6 {
>
> entry:
>
>   %call = tail call spir_func i128 @get_global_id(i32 0) #2
>
>   %conv = trunc i128 %call to i32
>
>   store i32 %conv, i32 addrspace(1)* %in0, align 4, !tbaa !7
>
>   ret void
>
> }
>
>
>
> In this case, get_global_id is not mangled to _Z13get_global_idj according
> to [2].
>
>
>
> Is there an option for clang or an LLVM clang to do the mangling for spir
> builtins?
>
>
>
> Thanks
>
> Hongbin
>
>
>
>
>
> [1] https://www.khronos.org/registry/spir/specs/spir_spec-2.0.pdf, page 36
>
> [2] https://github.com/KhronosGroup/SPIR-Tools/wiki/
> SPIR-2.0-built-in-functions
>
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20160912/1550f861/attachment.html>