[cfe-dev] builtins name mangling in SPIR 2.0

Mon Sep 19 03:53:12 PDT 2016

I'd like to clarify that old SPIR 1.2 and SPIR 2.0 formats are flavors of LLVM IR version 3.2/3.4, so clang 3.9 will never be able to produce "conformant" SPIR programs. Considering this I think it makes sense to evaluate the benefit of features from SPIR 1.2/2.0 specification we are trying to implement in ToT clang.

Could you clarify the problem you are trying to solve by mangling printf function?

Thanks,
Alexey

From: Anastasia Stulova [mailto:Anastasia.Stulova at arm.com]
Sent: Sunday, September 18, 2016 10:59 PM
To: Hongbin Zheng <etherzhhb at gmail.com>; Liu, Yaxun (Sam) <Yaxun.Liu at amd.com>
Cc: cfe-dev at lists.llvm.org; llvm-dev <llvm-dev at lists.llvm.org>; Bader, Alexey <alexey.bader at intel.com>; nd <nd at arm.com>
Subject: Re: 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<mailto:etherzhhb at gmail.com>>
Sent: 17 September 2016 05:32:54
To: Liu, Yaxun (Sam)
Cc: cfe-dev at lists.llvm.org<mailto:cfe-dev at lists.llvm.org>; llvm-dev; Bader, Alexey (alexey.bader at intel.com<mailto: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

--------------------------------------------------------------------
Joint Stock Company Intel A/O
Registered legal address: Krylatsky Hills Business Park,
17 Krylatskaya Str., Bldg 4, Moscow 121614,
Russian Federation

This e-mail and any attachments may contain confidential material for
the sole use of the intended recipient(s). Any review or distribution
by others is strictly prohibited. If you are not the intended
recipient, please contact the sender and delete all copies.
-------------- next part --------------
An HTML attachment was scrubbed...
URL: <http://lists.llvm.org/pipermail/cfe-dev/attachments/20160919/7e3aa77b/attachment.html>