extern __constant__ fails to link (global vars)

[PaulKarlshoeferBULL]

Dear developers,

I have "hipified" a large cuda code, using hipify clang. I run into an linker error and I lack a solution to this. Here is a minimal example:

fileA.h

#ifndef file_a_h
#define file_a_h
#include<hip/hip_runtime.h>

namespace my_names
{
    extern __constant__ int _MY_CONST;
}
#endif

fileA.cpp

#include <stdio.h>
#include "fileA.h"
#include <hip/hip_runtime.h>

namespace my_names
{
    __constant__ int _MY_CONST;
}

fileB.cpp

#include <stdio.h>
#include "fileA.h"

int main(int argc, char **argv){
    int tau = 1;
    hipError_t hip_err;
    hip_err = hipMemcpyToSymbol(HIP_SYMBOL(my_names::_MY_CONST), &tau, sizeof(tau), 0, hipMemcpyHostToDevice);

    return 0;
}

Compilation: /opt/rocm-4.0.0/bin/hipcc -c fileA.cpp -o fileA.o /opt/rocm-4.0.0/bin/hipcc -c fileB.cpp -o fileB.o /opt/rocm-4.0.0/bin/hipcc --hip-link fileA.o fileB.o -o test.exe

linker error: undefined reference to `my_names::_MY_CONST'

Question: How to get around this, preferably without changing the overall structure of the code. Without the constant modifier it would compile, but I want it to be on constant device memory. Thanks in advance!

[yxsamliu]

fileB.cpp is accessing a device-side variable defined in fileA.cpp. You need option -fgpu-rdc in both compilation and linking commands.

[PaulKarlshoeferBULL]

Thank you for your answer! I tried that already, like so (same error):

/opt/rocm-4.0.0/bin/hipcc -fgpu-rdc -c fileA.cpp -o fileA.o
/opt/rocm-4.0.0/bin/hipcc -fgpu-rdc -c fileB.cpp -o fileB.o
/opt/rocm-4.0.0/bin/hipcc --hip-link -fgpu-rdc fileA.o fileB.o -o test.exe

Output:

/tmp/fileB-298a1e.o: In function `main':
fileB.cpp:(.text+0xf): undefined reference to `my_names::_MY_CONST'
clang-12: error: linker command failed with exit code 1 (use -v to see invocation)

[yxsamliu]

There is a known issue about the linkage of shadow variable for -fgpu-rdc which has been fixed in clang trunk but the fix was not in ROCm release yet. https://reviews.llvm.org/D95901

As a workaround, you need to do hipMemcpyToSymbol in the same file as the device-side variable is defined, e.g. in fileB.cpp. You may define a function in fileB.cpp to do that and call it in fileA.cpp.

[PaulKarlshoeferBULL]

Thank you very much for pointing this out! It didn't occur to me that I might look at a bug:)

The workaround you propose works for the example above. I will change the code accordingly.