undefined protected symbol: enzyme_const when using HIP and -fgpu-rdc flag #2610
Description
When using HIP and the -fgpu-rdc flag, it seems like it's not possible to define device enzyme variables.
Steps to reproduce
- install enzyme with llvm-amdgpu (any version probably works)
- create the following main function
- build with the following bash command
- removing
-fgpu-rdcwill fix the problem - alternatively, removing the
DEVICEdefinition ofenzyme_constalso fixes the problem
Cpp
#include <iostream>
#define DEVICE __device__
#define HOST_DEVICE __host__ __device__
extern int enzyme_const;
extern int DEVICE enzyme_const; // comment this to additionally remove error
template <typename return_type, typename... Args>
HOST_DEVICE return_type __enzyme_autodiff(Args...);
double square(double x) {
return x * x;
}
void dsquare(double x) {
__enzyme_autodiff<double>((void*) square, enzyme_const, x);
}
int main() {
double x = 2;
dsquare(x);
#ifndef NOVERSION
std::cout << "Enzyme version: "
<< ENZYME_VERSION_MAJOR << " "
<< ENZYME_VERSION_MINOR << " "
<< ENZYME_VERSION_PATCH << "\n";
#endif
}Command
/opt/rocm-6.4.2/llvm/bin/amdclang++ main.hip \
--offload-arch=gfx942 \
-fplugin=/usr/workspace/meemee/smith/Enzyme/enzyme/build_toss4_cray/Enzyme/ClangEnzyme-19.so \
-fgpu-rdc
-o mainDemonstration of error
[meemee@rzadams1017:test_proj]$ /opt/rocm-6.4.2/llvm/bin/amdclang++ main.hip --offload-arch=gfx942 -fplugin=/usr/workspace/meemee/smith/Enzyme/enzyme/build_toss4_cray/Enzyme/ClangEnzyme-19.so -o main
[meemee@rzadams1017:test_proj]$ ./main
Enzyme version: 0 0 79
[meemee@rzadams1017:test_proj]$ /opt/rocm-6.4.2/llvm/bin/amdclang++ main.hip --offload-arch=gfx942 -fplugin=/usr/workspace/meemee/smith/Enzyme/enzyme/build_toss4_cray/Enzyme/ClangEnz
yme-19.so -o main -fgpu-rdc
lld: error: undefined protected symbol: enzyme_const
>>> referenced by /var/tmp/meemee/main-gfx942-1fe835.out.lto.o:(__clang_gpu_used_external)
clang++: error: amdgcn-link command failed with exit code 1 (use -v to see invocation)