[llvm-dev] Mangled SPIR function name

Jessica Clarke via llvm-dev llvm-dev at lists.llvm.org
Mon Jan 3 14:45:20 PST 2022


On 3 Jan 2022, at 21:33, Frank Winter via llvm-dev <llvm-dev at lists.llvm.org> wrote:
> 
> "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?
> 
> Frank

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.

Jess



More information about the llvm-dev mailing list