extern "C" (on declaration only) produces mangled symbols for __host__ functions

[jithunnair-amd]

Issue copied from: https://github.com/ROCm-Developer-Tools/HIP/issues/630 @whchung suggested filing HCC issue as it might be related to HCC Clang CodeGen.

If a __host__ function is declared as extern "C", but the definition isn't labelled as extern "C", hipcc/hcc produces a mangled symbol in the resulting object file. However, if the definition also has an extern "C" label, the symbol is non-mangled as expected. Also, this unexpected behaviour doesn't occur for __global__ functions (they produce the expected non-mangled symbol regardless of whether the definition also has an extern "C"). Below is a testcase for the same:

experiment.h: extern "C" void FuncA(int exp);

experiment.cu: #include <stdio.h> #include "experiment.h" #include "hip/hip_runtime.h"

/*extern "C"*/ __host__ void FuncA(int exp) { printf("I got %d", exp); }

Compile command: hipcc -fPIC -c experiment.cu -o experiment.o

rocm-user@51f7736e9394:~/experiment$ nm experiment.o | grep FuncA 00000000000038e0 T _Z5FuncAi

[ezyang]

This issue affected me today.

[srinivamd]

The presence (addition) of the __host__ macro/attribute in the function definition __host__ void FuncA(int exp)... might be the reason why the extern "C" declaration on FuncA() WITHOUT the __host__ attribute does not get applied. Does removing the __host__ attribute in the function definition make the demangling of FuncA work correctly? NOTE: That does not mean to suggest removing __host__ attribute but only to futher narrow the scope of this issue. Does __host__ attribute imply C++ linkage by default?

[ezyang]

IIRC, I tried the version where I declared __host__ extern "C" in header, and nothing in the cpp file, and I had the same problem. But this problem is really easy to repro, so you should just double check...

[srinivamd]

The header declarations are pretty much a no-op when the function definitions in the compiled translation unit is inconsistent - in this case the specifier __host__ makes them inconsistent and compiler seems to supersede the declaration. Plus, intervening inclusion on header files that may have different linkage might also be affecting I think. Granted that there's an issue, but FWIW here's modified version of your experiment code that can work: cat experiment.h #ifdef __cplusplus extern "C" { #endif __host__ void FuncA(int exp); #ifdef __cplusplus }; #endif

cat experiment.cu #include <stdio.h> #include "hip/hip_runtime.h" #include "experiment.h"

/*extern "C"*/ __host__ void FuncA(int exp) { printf("I got %d", exp); }

and hipcc -c experiment.cu nm experiment.o | grep Func 0000000000002ab0 T FuncA

HTH

[scchan]

I could only reproduce this issue if there's a mismatch between the function declaration and the definition. More specifically, when the declaration is missing the host keyword:

// failing extern "C" int foo(int a); __host__ int foo(int a) { return a; }

As a short term workaround, adding __host__ to the function declaration should make it work:

// working extern "C" __host__ int foo(int a); __host__ int foo(int a) { return a; }

Looking at the compiler generated AST, in the working case, it shows that there's a link (e.g. prev tag in the 2nd FunctionDecl)between the 2 different function declarations:

|-LinkageSpecDecl 0x55d972c721e8 <./cpu_extern_c.cpp:15:1, col:46> col:8 C | -FunctionDecl 0x55d972c72320 <col:12, col:46> col:37 foo 'int (int)' | |-ParmVarDecl 0x55d972c72250 <col:41, col:45> col:45 a 'int' | -CXXAMPRestrictCPUAttr 0x55d972c723c8 col:27 -FunctionDecl 0x55d972c72518 prev 0x55d972c72320 <line:16:1, line:18:1> line:16:26 foo 'int (int)' |-ParmVarDecl 0x55d972c72480 <col:30, col:34> col:34 used a 'int' |-CompoundStmt 0x55d972c72660 <col:37, line:18:1> | -ReturnStmt 0x55d972c72650 <line:17:3, col:10> | -ImplicitCastExpr 0x55d972c72638 <col:10> 'int' <LValueToRValue> | -DeclRefExpr 0x55d972c72618 col:10 'int' lvalue ParmVar 0x55d972c72480 'a' 'int' `-CXXAMPRestrictCPUAttr 0x55d972c725c0 line:16:16

In the failing case, the link is missing:

|-LinkageSpecDecl 0x56249fbe11c8 <./cpu_extern_c.cpp:10:1, col:25> col:8 C | -FunctionDecl 0x56249fbe1300 <col:12, col:25> col:16 foo 'int (int)' | -ParmVarDecl 0x56249fbe1230 <col:20, col:24> col:24 a 'int' -FunctionDecl 0x56249fbe14a0 <line:11:1, line:13:1> line:11:26 foo 'int (int)' |-ParmVarDecl 0x56249fbe1408 <col:30, col:34> col:34 used a 'int' |-CompoundStmt 0x56249fbe15e8 <col:37, line:13:1> | -ReturnStmt 0x56249fbe15d8 <line:12:3, col:10> | -ImplicitCastExpr 0x56249fbe15c0 <col:10> 'int' <LValueToRValue> | -DeclRefExpr 0x56249fbe15a0 col:10 'int' lvalue ParmVar 0x56249fbe1408 'a' 'int' `-CXXAMPRestrictCPUAttr 0x56249fbe1548 line:11:16

This indicates there's a signature matching bug in the early stages of the compiler frontend.