[llvm] r241081 - Force relocation mode to be default, regardless of what is passed to the backend.

[Samuel Antao]

Author: sfantao
Date: Tue Jun 30 12:18:00 2015
New Revision: 241081

URL: http://llvm.org/viewvc/llvm-project?rev=241081&view=rev
Log:
Force relocation mode to be default, regardless of what is passed to the backend.


Added:
    llvm/trunk/test/CodeGen/NVPTX/globals_lowering.ll
Modified:
    llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp

Modified: llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp?rev=241081&r1=241080&r2=241081&view=diff
==============================================================================
--- llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp (original)
+++ llvm/trunk/lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp Tue Jun 30 12:18:00 2015
@@ -54,7 +54,10 @@ createNVPTXMCSubtargetInfo(const Triple
 static MCCodeGenInfo *createNVPTXMCCodeGenInfo(
     StringRef TT, Reloc::Model RM, CodeModel::Model CM, CodeGenOpt::Level OL) {
   MCCodeGenInfo *X = new MCCodeGenInfo();
-  X->initMCCodeGenInfo(RM, CM, OL);
+
+  // The default relocation model is used regardless of what the client has
+  // specified, as it is the only relocation model currently supported.
+  X->initMCCodeGenInfo(Reloc::Default, CM, OL);
   return X;
 }
 

Added: llvm/trunk/test/CodeGen/NVPTX/globals_lowering.ll
URL: http://llvm.org/viewvc/llvm-project/llvm/trunk/test/CodeGen/NVPTX/globals_lowering.ll?rev=241081&view=auto
==============================================================================
--- llvm/trunk/test/CodeGen/NVPTX/globals_lowering.ll (added)
+++ llvm/trunk/test/CodeGen/NVPTX/globals_lowering.ll Tue Jun 30 12:18:00 2015
@@ -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
+}