[PATCH] Consider offset in global variables during lowering in NVPTX.

[Samuel Antao]

Hi jholewinski,

If the NVPTX backend is used with relocation model 'static', DAGCombiner will have global addresses combined with any add instruction that adds a constant offset to it. However, lowering is not prepared to deal with the offset that is added in GlobalAddress nodes. This issue can be replicated if llc is called with relocation-model=static or if the backend is invoked directly from clang.

This patch fixes the issue by undoing what DAGCombiner does, which looks to me as the less disruptive solution. Next to the fix, a comment is added explaining other alternatives to tackle this issue. 

Thanks!
Samuel

http://reviews.llvm.org/D9940

Files:
  lib/Target/NVPTX/NVPTXISelLowering.cpp
  test/CodeGen/NVPTX/globals_lowering.ll

Index: lib/Target/NVPTX/NVPTXISelLowering.cpp
===================================================================
--- lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -877,9 +877,24 @@
 SDValue
 NVPTXTargetLowering::LowerGlobalAddress(SDValue Op, SelectionDAG &DAG) const {
   SDLoc dl(Op);
-  const GlobalValue *GV = cast<GlobalAddressSDNode>(Op)->getGlobal();
+
+  GlobalAddressSDNode *GN = cast<GlobalAddressSDNode>(Op);
+
+  const GlobalValue *GV = GN->getGlobal();
   Op = DAG.getTargetGlobalAddress(GV, dl, getPointerTy());
-  return DAG.getNode(NVPTXISD::Wrapper, dl, getPointerTy(), Op);
+  Op = DAG.getNode(NVPTXISD::Wrapper, dl, getPointerTy(), Op);
+
+  // We need to consider any offset that comes with the global. Basically, here
+  // we undo what the DAG combiner does when the relocation model is set to
+  // static. We could alternatively set the relocation model to default in
+  // CodeGenInfo or make isOffsetFoldingLegal() from TargetLowering to depend on
+  // some target dependent hook that could be overloaded for NVPTX.
+  if (GN->getOffset()) {
+    SDValue Offset = DAG.getConstant(GN->getOffset(), dl, getPointerTy());
+    Op = DAG.getNode(ISD::ADD, Op, getPointerTy(), Op, Offset);
+  }
+
+  return Op;
 }
 
 std::string
Index: test/CodeGen/NVPTX/globals_lowering.ll
===================================================================
--- /dev/null
+++ test/CodeGen/NVPTX/globals_lowering.ll
@@ -0,0 +1,15 @@
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -relocation-model=static | FileCheck %s --check-prefix CHK
+
+%MyStruct = type { i32, i32, float }
+ at Gbl = internal addrspace(3) global [1024 x %MyStruct] zeroinitializer
+
+; CHK-LABEL: foo
+define void @foo(float %f) {
+entry:
+  ; CHK: ld.shared.f32  %{{[a-zA-Z0-9]+}}, [Gbl+8];
+  %0 = load float, float addrspace(3)* getelementptr inbounds ([1024 x %MyStruct], [1024 x %MyStruct] addrspace(3)* @Gbl, i32 0, i32 0, i32 2)
+  %add = fadd float %0, %f
+  ; CHK: st.shared.f32   [Gbl+8], %{{[a-zA-Z0-9]+}};
+  store float %add, float addrspace(3)* getelementptr inbounds ([1024 x %MyStruct], [1024 x %MyStruct] addrspace(3)* @Gbl, i32 0, i32 0, i32 2)
+  ret void
+}

EMAIL PREFERENCES
  http://reviews.llvm.org/settings/panel/emailpreferences/
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D9940.26331.patch
Type: text/x-patch
Size: 2187 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20150522/af076c17/attachment.bin>