diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
--- a/clang/lib/Driver/ToolChains/Clang.cpp
+++ b/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
diff --git a/clang/test/CodeGen/amdgpu-alias-undef-symbols.hip b/clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
new file mode 100644
--- /dev/null
+++ b/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) {
+
+}
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
--- a/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUTargetMachine.cpp
@@ -689,8 +689,6 @@
         if (InternalizeSymbols) {
           PM.addPass(GlobalDCEPass());
         }
-        if (EarlyInlineAll && !EnableFunctionCalls)
-          PM.addPass(AMDGPUAlwaysInlinePass());
       });
 
   PB.registerCGSCCOptimizerLateEPCallback(