Mangled SPIR function name

“clang -triple spir64 …” would generate from an OpenCL statement like

uint idx = get_global_id(0);

something like this:

%0 = call i64 @_Z13get_global_idj(i32 0)

Note it generates the mangled form.

Using only LLVM builder tools, i.e. no clang, how can I generate this mangled form?

Using Module::getOrInsertFunction( “_Z13get_global_idj” , … ) seems little portable.

Assuming that I’m building a “spir_kernel”, is the above (simple) solution safe enough?


This is just the Itanium C++ ABI’s mangling, implemented in
clang/lib/AST/ItaniumMangle.cpp. This is because all the OpenCL
builtins are defined as overloadable functions and Clang implements
overloadable functions by mangling them as if they were C++ functions
(since the overloadable attribute just asks Clang to use C++’s
overloading semantics), which is Itanium everywhere except for
Windows’s ABI, which does its own thing.

_Z is the prefix for all such mangled symbols. 13 is the length of
get_global_id. j indicates that it takes an unsigned int argument. If
the argument is always an unsigned int and never some other primitive
(e.g. unsigned long) then it should be portable enough, and according
to clang/lib/Headers/opencl-c.h it is always an unsigned int argument.