Consider offset in global variables during lowering in NVPTX.

[sfantao]

Summary

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

[sfantao]

Ping!

[sfantao]

Ping!

[jholewinski]

What is the purpose to using a static relocation model for NVPTX? I would prefer to just enforce a default here, especially since this is essentially unused for NVPTX.

[sfantao]

This diff changes the original fix to enforce the default relocation model to be used regardless of what the client specifies.

[sfantao]

My understanding is that the backend clients (e.g. clang) may request a given relocation model. If nothing else is specified by the user, this relocation model would be 'static'. I thought relocation model = default would map to one of the supported relocation models, and static would be what makes more sense for NVPTX. This would however drive some optimized behaviors in the combiner causing the issue reported here to be exposed.

The diff above discards any relocation model information specified by the client.

Let me know any other comments you may have,

Thanks!
Samuel

[sfantao]

Any more comments on this patch?

Thanks!
Samuel

[sfantao]

Any other comments about this patch?

Thanks,
Samuel

[sfantao]

Ping!

[jholewinski]

Sounds reasonable to me.

[sfantao]

Committed in r241081.

Thanks!
Samuel

lib/Target/NVPTX/MCTargetDesc/NVPTXMCTargetDesc.cpp

createNVPTXMCSubtargetInfo(StringRef TT, StringRef CPU, StringRef FS) {
   MCSubtargetInfo *X = new MCSubtargetInfo();
   InitNVPTXMCSubtargetInfo(X, TT, CPU, FS);
   return X;
 }
 
 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;
 }
 
 static MCInstPrinter *createNVPTXMCInstPrinter(const Triple &T,
                                                unsigned SyntaxVariant,
                                                const MCAsmInfo &MAI,
                                                const MCInstrInfo &MII,
                                                const MCRegisterInfo &MRI) {

test/CodeGen/NVPTX/globals_lowering.ll

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