[PATCH] D109707: [HIP] [AlwaysInliner] Disable AlwaysInliner to eliminate undefined symbols

[Anshil Gandhi via Phabricator via cfe-commits]

gandhi21299 updated this revision to Diff 372551.
gandhi21299 added a comment.

- added the include header for HIP runtime


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D109707/new/

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
@@ -691,7 +691,7 @@
           PM.addPass(GlobalDCEPass());
         }
         if (EarlyInlineAll && !EnableFunctionCalls)
-          PM.addPass(AMDGPUAlwaysInlinePass());
+          PM.addPass(AMDGPUAlwaysInlinePass(false));
       });
 
   PB.registerCGSCCOptimizerLateEPCallback(
Index: clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
@@ -0,0 +1,16 @@
+// RUN: %clang --offload-arch=gfx906 --cuda-device-only -emit-llvm -S -o - %s \
+// RUN:   -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm 
-amdgpu-function-calls=false | \
+// RUN:   FileCheck %s
+
+#include "hip/hip_runtime.h"
+
+// 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
@@ -5084,9 +5084,9 @@
   }
 
   // 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())
+  // a linker bug (see <rdar://problem/7651567>), and CUDA device code, where
+  // aliases aren't supported.
+  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
@@ -691,7 +691,7 @@
           PM.addPass(GlobalDCEPass());
         }
         if (EarlyInlineAll && !EnableFunctionCalls)
-          PM.addPass(AMDGPUAlwaysInlinePass());
+          PM.addPass(AMDGPUAlwaysInlinePass(false));
       });
 
   PB.registerCGSCCOptimizerLateEPCallback(
Index: clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
===================================================================
--- /dev/null
+++ clang/test/CodeGen/amdgpu-alias-undef-symbols.hip
@@ -0,0 +1,16 @@
+// RUN: %clang --offload-arch=gfx906 --cuda-device-only -emit-llvm -S -o - %s \
+// RUN:   -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false | \
+// RUN:   FileCheck %s
+
+#include "hip/hip_runtime.h"
+
+// 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
@@ -5084,9 +5084,9 @@
   }
 
   // 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())
+  // a linker bug (see <rdar://problem/7651567>), and CUDA device code, where
+  // aliases aren't supported.
+  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