[cfe-dev] builtins name mangling in SPIR 2.0

[Hongbin Zheng via cfe-dev]

Hi Sam and others,

I saw that in [1], printf is mangled to _Z6printfPrU3AS2cz, while in
clang's opencl-c.h[2], printf does not have the overload attribute:
int printf(__constant const char* st, ...); (and it is different from the
standard, which is printf(restrict __constant char *, ...))
I try the following code:
#include <opencl-c.h>

__kernel void vadd(__global const int* a, __global const int* b, __global
int* c) {
  printf("aaaaa");
}

and get a printf that is not mangled:
  %call = tail call spir_func i32 (i8 addrspace(2)*, ...) @printf(i8
addrspace(2)* getelementptr inbounds ([6 x i8], [6 x i8] addrspace(2)*
@.str, i64 0, i64 0))

with the following command line:
clang -cc1 -internal-isystem
/wrk/xsjhdnobkup2/hongbinz/omp/build-llvm/bin/../lib/clang/3.9.0/include
-nostdsysteminc -S -emit-llvm -o -

Is this the correct behavior?

Thanks
Hongbin

[1]
https://github.com/KhronosGroup/SPIR-Tools/wiki/SPIR-2.0-built-in-functions
[2] http://clang.llvm.org/doxygen/opencl-c_8h_source.html

On Mon, Sep 12, 2016 at 1:43 PM, Hongbin Zheng <etherzhhb at gmail.com> wrote:

> 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/20160915/e5c71c9e/attachment.html>