builtins name mangling in SPIR 2.0

Project Infrastructure LLVM Dev List Archives
1

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

2

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

3

Thanks a lot.

4

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

5
  • 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

6

+ 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

7

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

8

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?

9

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?

I am not solving any problem by mangling the printf. I just confused when I
saw the header and the SPIR 2.0 "Provisional" specs do not agree :slight_smile: