Liu, Yaxun (Sam) via llvm-dev
2016-Sep-16 15:10 UTC
[llvm-dev] builtins name mangling in SPIR 2.0
+ Alexey Anastasia
According to SPIR spec v1.2 s2.10.3
2.10.3 The printf function
The printf function is supported, and is mangled according to its prototype as
follows:
int printf(constant char * restrict fmt, ... )
Note that the ellipsis formal argument (...) is mangled to argument type
specifier z
It seems printf should be mangled.
Alexey/Anastasia,
What do you think? Thanks.
Sam
From: Hongbin Zheng [mailto:etherzhhb at gmail.com]
Sent: Friday, September 16, 2016 1:54 AM
To: Liu, Yaxun (Sam) <Yaxun.Liu at amd.com>
Cc: cfe-dev at lists.llvm.org; llvm-dev <llvm-dev at lists.llvm.org>
Subject: Re: builtins name mangling in SPIR 2.0
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<mailto: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<mailto: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<mailto:etherzhhb at
gmail.com>]
Sent: Monday, September 12, 2016 4:21 PM
To: cfe-dev at lists.llvm.org<mailto:cfe-dev at lists.llvm.org>; llvm-dev
<llvm-dev at lists.llvm.org<mailto:llvm-dev at lists.llvm.org>>;
Liu, Yaxun (Sam) <Yaxun.Liu at amd.com<mailto: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/llvm-dev/attachments/20160916/30d42605/attachment.html>
Hongbin Zheng via llvm-dev
2016-Sep-17 04:32 UTC
[llvm-dev] builtins name mangling in SPIR 2.0
On Fri, Sep 16, 2016 at 8:10 AM, Liu, Yaxun (Sam) <Yaxun.Liu at amd.com> wrote:> + Alexey Anastasia > > > > According to SPIR spec v1.2 s2.10.3 > > > > *2.10.3 The printf function* > The printf function is supported, and is mangled according to its > prototype as follows: > int printf(constant char * restrict fmt, ... ) >clang 3.9 will drop the 'restrict' when it mangle this name, or I am missing some commandline argument? Thanks Hongbin> Note that the ellipsis formal argument (...) is mangled to argument type > specifier z > > > > It seems printf should be mangled. > > > > Alexey/Anastasia, > > > > What do you think? Thanks. > > > > Sam > > > > *From:* Hongbin Zheng [mailto:etherzhhb at gmail.com] > *Sent:* Friday, September 16, 2016 1:54 AM > *To:* Liu, Yaxun (Sam) <Yaxun.Liu at amd.com> > *Cc:* cfe-dev at lists.llvm.org; llvm-dev <llvm-dev at lists.llvm.org> > *Subject:* Re: builtins name mangling in SPIR 2.0 > > > > 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/llvm-dev/attachments/20160916/21ef8d9d/attachment.html>
Anastasia Stulova via llvm-dev
2016-Sep-18 19:59 UTC
[llvm-dev] builtins name mangling in SPIR 2.0
I don't see any problem mangling it to be honest even though there seems to
be only one prototype anyways.
We could add restrict in as well.
Cheers,
Anastasia
________________________________
From: Hongbin Zheng <etherzhhb at gmail.com>
Sent: 17 September 2016 05:32:54
To: Liu, Yaxun (Sam)
Cc: cfe-dev at lists.llvm.org; llvm-dev; Bader, Alexey (alexey.bader at
intel.com); Anastasia Stulova
Subject: Re: builtins name mangling in SPIR 2.0
On Fri, Sep 16, 2016 at 8:10 AM, Liu, Yaxun (Sam) <Yaxun.Liu at
amd.com<mailto:Yaxun.Liu at amd.com>> wrote:
+ Alexey Anastasia
According to SPIR spec v1.2 s2.10.3
2.10.3 The printf function
The printf function is supported, and is mangled according to its prototype as
follows:
int printf(constant char * restrict fmt, ... )
clang 3.9 will drop the 'restrict' when it mangle this name, or I am
missing some commandline argument?
Thanks
Hongbin
Note that the ellipsis formal argument (...) is mangled to argument type
specifier z
It seems printf should be mangled.
Alexey/Anastasia,
What do you think? Thanks.
Sam
From: Hongbin Zheng [mailto:etherzhhb at gmail.com<mailto:etherzhhb at
gmail.com>]
Sent: Friday, September 16, 2016 1:54 AM
To: Liu, Yaxun (Sam) <Yaxun.Liu at amd.com<mailto:Yaxun.Liu at
amd.com>>
Cc: cfe-dev at lists.llvm.org<mailto:cfe-dev at lists.llvm.org>; llvm-dev
<llvm-dev at lists.llvm.org<mailto:llvm-dev at lists.llvm.org>>
Subject: Re: builtins name mangling in SPIR 2.0
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<mailto: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<mailto: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<mailto:etherzhhb at
gmail.com>]
Sent: Monday, September 12, 2016 4:21 PM
To: cfe-dev at lists.llvm.org<mailto:cfe-dev at lists.llvm.org>; llvm-dev
<llvm-dev at lists.llvm.org<mailto:llvm-dev at lists.llvm.org>>;
Liu, Yaxun (Sam) <Yaxun.Liu at amd.com<mailto: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/llvm-dev/attachments/20160918/229b4d4e/attachment.html>