lld: error: undefined protected symbol: cCriticality >>> referenced by /tmp/FinalizeGPU-e78412/FinalizeGPU-gfx928.o:(FinalizeKeff_GPUg())

[DNYZH]

Environment

• HIP version: /rocm/dtk/25.04.1

Issue Description

I am porting a CUDA-based program to HIP using the hipify tool. After conversion, compilation fails with "undefined protected symbol" errors for multiple variables/functions (e.g., cCriticality). The error occurs during linking:

lld: error: undefined protected symbol: cCriticality  
>>> referenced by /tmp/FinalizeGPU-e78412/FinalizeGPU-gfx928.o:(FinalizeKeff_GPUg())  

Code Context

1. Variable Definition:

   • In .cpp file: Ddecl_devVar(cCriticality, CDCriticality_GPU)
   • In .h file: Dextern_devVar(cCriticality, CDCriticality_GPU)
     Many other symbols and functions exhibit the same issue.

2. CMake Configuration (GPU part):

file(GLOB SOURCES_HOST src/GPU/*.cpp)
file(GLOB HEADERS_HOST src/GPU/include/*.h)
file(GLOB SOURCES_CUDA src/GPU/CUDA/*.cpp)
file(GLOB HEADERS_CUDA src/GPU/CUDA_include/*.h)

add_custom_target(include_dep DEPENDS ${HEADERS_HOST} ${HEADERS_CUDA})

set_source_files_properties(${SOURCES_CUDA} ${SOURCES_HOST} PROPERTIES HIP_SOURCE_PROPERTY_FORMAT 1)

set(RMC_GPU "RMC_GPU")
hip_add_library(${RMC_GPU} STATIC ${SOURCES_CUDA} ${SOURCES_HOST} ${HEADERS_HOST} ${HEADERS_CUDA})
add_dependencies(${RMC_GPU} include_dep)

set_target_properties(${RMC_GPU} PROPERTIES HIP_ARCHITECTURES gfx906)

target_compile_options(${RMC_GPU} PUBLIC 
    $<$<COMPILE_LANGUAGE:CXX>:-fgpu-rdc>
    $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>
    -g -gdwarf-4
)

target_link_options(${RMC_GPU} PUBLIC
  "$<$<LINK_LANGUAGE:CXX>:-fgpu-rdc;--hip-link>"
  "$<$<LINK_LANGUAGE:HIP>:-fgpu-rdc;--hip-link>"
)

target_include_directories(${RMC_GPU} PUBLIC ... )  # (Abbreviated for clarity)

Request for Help

How should HIP device variables be declared/exported to resolve these linker errors? Thanks for your consideration!

[amd-nicknick]

Hi @DNYZH, I might need further context on this. Could you share partial reproducible CUDA code that is having issue converting? Let's align our configuration first.

[DNYZH]

CUDA Program Implementation

Thanks for your attention. In a '.h' file ,we define:

class CDCriticality_GPU {
public:
    CDCriticality_GPU() = default;

    int p_nNeuNumPerCyc;
    int p_nTotCycNum;
    int p_nInactCycNum;
    int p_nTotActCyc;

    bool autoConvergeByWd;
    bool p_bUseBatch;
    int p_nBatchNum;
    int p_nBatchLen;
    int p_nCycleNeuNum;
    int p_nCurrActiveCYCLE;
    int p_nCurrentCYCLE;
    int p_nCurrentPARTICLE;
    double p_dKeffFinal;
    
    enum p_eEstimate {
        Collision = 0,
        Absorption,
        TrackLength = 2,
        Covw
    };
    
    double p_dKeffWgtSum[3];
    double p_dKeffCyc[3];
    double p_dKeffSum[3];
    double p_dKeffCrosSum[3][3];
    double p_dKeffIndAve[3];
    double p_dKeffIndStd[3];
    double p_dKeffCovwAve[4];
    double p_dKeffCovwStd[4];

    double p_dTotStartWgt;
    long long p_llTotCollisionCount;
    double p_dStartWgt;
    int p_nMissParCount;
    int p_nMaxMissParCount;

    bool p_bMGAdjFisMatrix;

    __device__ void EstimateKeffCol_GPUd(const double dWgt, const double dMacroMuFisXs, const double dMacroTotXs);
    __device__ void EstimateKeffAbs_GPUd(const double dWgt, const double dNu, const double dMicroFisXs, const double dMicroTotXs);
    __device__ void EstimateKeffTrk_GPUd(const double dWgt, const double dMacroMuFisXs, const double dTrackLen);
    
    CDCriticality_GPU& operator=(const CDCriticality_t& other);

    bool p_bIsActiveCyc;
};

Variable Definition in .cu File

Ddecl_devVar(cCriticality, CDCriticality_GPU)

External Declaration in .h File

Dextern_devVar(cCriticality, CDCriticality_GPU)

Variable Usage in .cu File

__global__ void FinalizeKeff_GPUg() {
    printf("\nFinal Keff: %f      Standard Deviation: %f\n", 
           cCriticality.p_dKeffCovwAve[3],
           cCriticality.p_dKeffCovwStd[3]);
}

The error mentioned above then occurred. It should be noted that while using the hipify tool for conversion, the process proceeded without errors, but the converted files encountered issues during compilation. And in CUDA, there is no error when compiling. In addition, we attempted to combine all .cpp files into a single .cpp file for compilation, which resolves the aforementioned errors. Please contact me promptly if you need more information!

[amd-nicknick]

Sorry for the delayed response, I was caught up in urgent internal issue. With your repro steps, I tested two scenarios:

Common code files:

define.cu.hip (Hipify)

#include <hip/hip_runtime.h>
#include "Criticality.h"

__device__ CDCriticality_GPU cCriticality = {10};

main.cu.hip (Hipify)

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

__global__ void FinalizeKeff_GPUg(int *out) {
    printf("\nFinal Keff: %d      Standard Deviation: %d\n",
           cCriticality.p_nNeuNumPerCyc,
           cCriticality.p_nNeuNumPerCyc);
    out[0] = cCriticality.p_nNeuNumPerCyc;
}

int main() {
    int *out;
    int code = hipMallocManaged(&out, sizeof(int));
    printf("hipMalloc %d\n", code);
    FinalizeKeff_GPUg<<<1,1>>>(out);
    printf("Kern %d\n", hipDeviceSynchronize());
    printf("Bye %d\n", out[0]);
    return 0;
}

Criticality.h

class CDCriticality_GPU {
public:
    CDCriticality_GPU() = default;

    int p_nNeuNumPerCyc;
    int p_nTotCycNum;
    int p_nInactCycNum;
    int p_nTotActCyc;

    bool autoConvergeByWd;
    bool p_bUseBatch;
    int p_nBatchNum;
    int p_nBatchLen;
    int p_nCycleNeuNum;
    int p_nCurrActiveCYCLE;
    int p_nCurrentCYCLE;
    int p_nCurrentPARTICLE;
    double p_dKeffFinal;

    enum p_eEstimate {
        Collision = 0,
        Absorption,
        TrackLength = 2,
        Covw
    };

    double p_dKeffWgtSum[3];
    double p_dKeffCyc[3];
    double p_dKeffSum[3];
    double p_dKeffCrosSum[3][3];
    double p_dKeffIndAve[3];
    double p_dKeffIndStd[3];
    double p_dKeffCovwAve[4];
    double p_dKeffCovwStd[4];

    double p_dTotStartWgt;
    long long p_llTotCollisionCount;
    double p_dStartWgt;
    int p_nMissParCount;
    int p_nMaxMissParCount;

    bool p_bMGAdjFisMatrix;

    //__device__ void EstimateKeffCol_GPUd(const double dWgt, const double dMacroMuFisXs, const double dMacroTotXs);
    //__device__ void EstimateKeffAbs_GPUd(const double dWgt, const double dNu, const double dMicroFisXs, const double dMicroTotXs);
    //__device__ void EstimateKeffTrk_GPUd(const double dWgt, const double dMacroMuFisXs, const double dTrackLen);

    //CDCriticality_GPU& operator=(const CDCriticality_t& other);

    bool p_bIsActiveCyc;
};

__device__ extern CDCriticality_GPU cCriticality;

1. Single target

CMake:

add_executable(${ProjectId} define.cu.hip main.cu.hip)
set_source_files_properties(PROPERTIES LANGUAGE HIP)

target_link_libraries(${ProjectId} PRIVATE hip::device)

target_compile_options(${ProjectId} PUBLIC 
    $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>
)

target_link_options(${ProjectId} PUBLIC
    $<$<LINK_LANGUAGE:HIP>:-fgpu-rdc>
)

2. Static library target

CMake:

set_source_files_properties(PROPERTIES LANGUAGE HIP)

add_executable(${ProjectId} main.cu.hip)

add_library(${ProjectId}_lib define.cu.hip)

target_compile_options(${ProjectId} PUBLIC 
    $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>
)

target_compile_options(${ProjectId}_lib PUBLIC 
    $<$<COMPILE_LANGUAGE:HIP>:-fgpu-rdc>
)

target_link_options(${ProjectId} PUBLIC
    $<$<LINK_LANGUAGE:HIP>:-fgpu-rdc>
)

target_link_libraries(${ProjectId} 
    PRIVATE hip::device
    PRIVATE ${ProjectId}_lib)

Both of these cases compile and run correctly.

The error you mentioned:

lld: error: undefined protected symbol: cCriticality  
>> referenced by /tmp/FinalizeGPU-e78412/FinalizeGPU-gfx928.o:(FinalizeKeff_GPUg())  

Is indicative to either incorrect declaration of extern, or the -fgpu-rdc flag missing. The RDC option must be enabled for your scenario and must be present in both compilation and linking phase. Try adding --verbose to your CMake build command, this will dump out the actual compile and link commands. Ensure -fgpu-rdc is present in all invocation when HIP device / global code is compiled.

Thanks!

[amd-nicknick]

Hi @DNYZH, I'm closing this issue for now. Feel free to reopen or create a new issue should you encounter any further issues. Thanks!

[DNYZH]

Thank you very much for your help. I apologize for the delayed response. In fact, the issue stems from another variable not being properly defined. Due to the absence of a virtual destructor, the file containing the definition of this variable failed to compile. As a result, other variables—including cCriticality—were also not compiled correctly. We are currently working on resolving the subsequent issues. Thank you again for your great assistance!