[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