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 }
+@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
+}