gandhi21299 created this revision.
gandhi21299 added reviewers: yaxunl, aeubanks.
Herald added subscribers: foad, kerbowa, hiraditya, tpr, nhaehnle, jvesely,
arsenm.
gandhi21299 requested review of this revision.
Herald added projects: clang, LLVM.
Herald added subscribers: llvm-commits, cfe-commits.
By default clang emits complete contructors as alias of base constructors if
they are the same. The backend is supposed to emit symbols for the alias,
otherwise it causes undefined symbols. @yaxunl observed that this issue is
related to the llvm options `-amdgpu-early-inline-all=true` and
`-amdgpu-function-calls=false`. Disabling the AlwaysInliner fixes this issue.
Repository:
rG LLVM Github Monorepo
https://reviews.llvm.org/D109707
Files:
clang/lib/Driver/ToolChains/Clang.cpp
clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
Index: llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -689,8 +689,6 @@
if (InternalizeSymbols) {
PM.addPass(GlobalDCEPass());
}
- if (EarlyInlineAll && !EnableFunctionCalls)
- PM.addPass(AMDGPUAlwaysInlinePass());
});
PB.registerCGSCCOptimizerLateEPCallback(
Index: clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
@@ -0,0 +1,14 @@
+// RUN: %clang -c --offload-arch=gfx906 --cuda-device-only -emit-llvm -S -o -
%s
+// -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm
-amdgpu-function-calls=false
+// FileCheck %s
+
+// CHECK: %struct.B = type { i8 }
+struct B {
+
+ // CHECK: @_ZN1BC1Ei = hidden unnamed_addr alias void (%struct.B*, i32),
void (%struct.B*, i32)* @_ZN1BC2Ei
+ __device__ B(int x);
+};
+
+__device__ B::B(int x) {
+
+}
Index: clang/lib/Driver/ToolChains/Clang.cpp
===================================================================
--- clang/lib/Driver/ToolChains/Clang.cpp
+++ clang/lib/Driver/ToolChains/Clang.cpp
@@ -5066,7 +5066,7 @@
// Enable -mconstructor-aliases except on darwin, where we have to work
around
// a linker bug (see <rdar://problem/7651567>), and CUDA/AMDGPU device code,
// where aliases aren't supported.
- if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX() && !RawTriple.isAMDGPU())
+ if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX())
CmdArgs.push_back("-mconstructor-aliases");
// Darwin's kernel doesn't support guard variables; just die if we
Index: llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -689,8 +689,6 @@
if (InternalizeSymbols) {
PM.addPass(GlobalDCEPass());
}
- if (EarlyInlineAll && !EnableFunctionCalls)
- PM.addPass(AMDGPUAlwaysInlinePass());
});
PB.registerCGSCCOptimizerLateEPCallback(
Index: clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
@@ -0,0 +1,14 @@
+// RUN: %clang -c --offload-arch=gfx906 --cuda-device-only -emit-llvm -S -o - %s
+// -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false
+// FileCheck %s
+
+// CHECK: %struct.B = type { i8 }
+struct B {
+
+ // CHECK: @_ZN1BC1Ei = hidden unnamed_addr alias void (%struct.B*, i32), void (%struct.B*, i32)* @_ZN1BC2Ei
+ __device__ B(int x);
+};
+
+__device__ B::B(int x) {
+
+}
Index: clang/lib/Driver/ToolChains/Clang.cpp
===================================================================
--- clang/lib/Driver/ToolChains/Clang.cpp
+++ clang/lib/Driver/ToolChains/Clang.cpp
@@ -5066,7 +5066,7 @@
// Enable -mconstructor-aliases except on darwin, where we have to work around
// a linker bug (see <rdar://problem/7651567>), and CUDA/AMDGPU device code,
// where aliases aren't supported.
- if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX() && !RawTriple.isAMDGPU())
+ if (!RawTriple.isOSDarwin() && !RawTriple.isNVPTX())
CmdArgs.push_back("-mconstructor-aliases");
// Darwin's kernel doesn't support guard variables; just die if we
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
