gandhi21299 updated this revision to Diff 376559.
gandhi21299 added a comment.
- Since callees may alias to a function pointer, it makes sense for
`getCalleeFunction(...)` to return a `Function` which is a cast of the operand
of a `GlobalAlias`.
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/CodeGenCUDA/amdgpu-alias-undef-symbols.cu
llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp
llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp
llvm/lib/Target/AMDGPU/SIISelLowering.cpp
llvm/test/CodeGen/AMDGPU/inline-calls.ll
Index: llvm/test/CodeGen/AMDGPU/inline-calls.ll
===================================================================
--- llvm/test/CodeGen/AMDGPU/inline-calls.ll
+++ llvm/test/CodeGen/AMDGPU/inline-calls.ll
@@ -1,6 +1,5 @@
; RUN: llc -march=amdgcn -mcpu=tahiti -verify-machineinstrs < %s | FileCheck %s
; RUN: llc -march=amdgcn -mcpu=tonga -verify-machineinstrs < %s | FileCheck %s
-; RUN: llc -march=r600 -mcpu=redwood -verify-machineinstrs < %s | FileCheck %s
; ALL-NOT: {{^}}func:
define internal i32 @func(i32 %a) {
@@ -18,8 +17,8 @@
ret void
}
-; CHECK-NOT: func_alias
-; ALL-NOT: func_alias
+; CHECK: func_alias
+; ALL: func_alias
@func_alias = alias i32 (i32), i32 (i32)* @func
; ALL: {{^}}kernel3:
Index: llvm/lib/Target/AMDGPU/SIISelLowering.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/SIISelLowering.cpp
+++ llvm/lib/Target/AMDGPU/SIISelLowering.cpp
@@ -3007,6 +3007,7 @@
bool IsSibCall = false;
bool IsThisReturn = false;
MachineFunction &MF = DAG.getMachineFunction();
+ GlobalAddressSDNode *GSD = dyn_cast<GlobalAddressSDNode>(Callee);
if (Callee.isUndef() || isNullConstant(Callee)) {
if (!CLI.IsTailCall) {
@@ -3264,7 +3265,7 @@
Ops.push_back(Callee);
// Add a redundant copy of the callee global which will not be legalized, as
// we need direct access to the callee later.
- if (GlobalAddressSDNode *GSD = dyn_cast<GlobalAddressSDNode>(Callee)) {
+ if (GSD) {
const GlobalValue *GV = GSD->getGlobal();
Ops.push_back(DAG.getTargetGlobalAddress(GV, DL, MVT::i64));
} else {
Index: llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUResourceUsageAnalysis.cpp
@@ -29,6 +29,8 @@
#include "SIMachineFunctionInfo.h"
#include "llvm/Analysis/CallGraph.h"
#include "llvm/CodeGen/TargetPassConfig.h"
+#include "llvm/IR/GlobalAlias.h"
+#include "llvm/IR/GlobalValue.h"
#include "llvm/Target/TargetMachine.h"
using namespace llvm;
@@ -61,7 +63,8 @@
assert(Op.getImm() == 0);
return nullptr;
}
-
+ if (auto *GA = dyn_cast<GlobalAlias>(Op.getGlobal()))
+ return cast<Function>(GA->getOperand(0));
return cast<Function>(Op.getGlobal());
}
Index: llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUCallLowering.cpp
@@ -913,14 +913,17 @@
if (Info.Callee.isReg()) {
CallInst.addReg(Info.Callee.getReg());
CallInst.addImm(0);
- } else if (Info.Callee.isGlobal() && Info.Callee.getOffset() == 0) {
- // The call lowering lightly assumed we can directly encode a call target in
- // the instruction, which is not the case. Materialize the address here.
+ } else if (Info.Callee.isGlobal()) {
const GlobalValue *GV = Info.Callee.getGlobal();
- auto Ptr = MIRBuilder.buildGlobalValue(
- LLT::pointer(GV->getAddressSpace(), 64), GV);
- CallInst.addReg(Ptr.getReg(0));
- CallInst.add(Info.Callee);
+ if (Info.Callee.getOffset() == 0) {
+ // The call lowering lightly assumed we can directly encode a call target
+ // in the instruction, which is not the case. Materialize the address
+ // here.
+ auto Ptr = MIRBuilder.buildGlobalValue(
+ LLT::pointer(GV->getAddressSpace(), 64), GV);
+ CallInst.addReg(Ptr.getReg(0));
+ CallInst.add(Info.Callee);
+ }
} else
return false;
Index: llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp
===================================================================
--- llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp
+++ llvm/lib/Target/AMDGPU/AMDGPUAlwaysInlinePass.cpp
@@ -93,6 +93,8 @@
for (GlobalAlias &A : M.aliases()) {
if (Function* F = dyn_cast<Function>(A.getAliasee())) {
+ if (A.getLinkage() != GlobalValue::InternalLinkage)
+ continue;
A.replaceAllUsesWith(F);
AliasesToRemove.push_back(&A);
}
Index: clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu
===================================================================
--- /dev/null
+++ clang/test/CodeGenCUDA/amdgpu-alias-undef-symbols.cu
@@ -0,0 +1,17 @@
+// REQUIRES: amdgpu-registered-target, clang-driver
+
+// RUN: %clang --offload-arch=gfx906 --cuda-device-only -x hip -emit-llvm -S -o - %s \
+// RUN: -fgpu-rdc -O3 -mllvm -amdgpu-early-inline-all=true -mllvm -amdgpu-function-calls=false | \
+// RUN: FileCheck %s
+
+#include "Inputs/cuda.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
@@ -5102,9 +5102,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
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
