[NVPTX] Lower loads from global constants using ld.global.nc (aka LDG).

[jlebar]

Summary

After D43914, loads from global variables in addrspace(1) happen with
ld.global. But since they're constants, even better would be to use
ld.global.nc, aka ldg.

[tra]

Test nit. LGTM otherwise.

llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll

19

Should there be a CHECK: ld.global.nc.... here and below?

[jlebar]

llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll

19

Haha, oops. :) Thanks.

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

 static bool canLowerToLDG(MemSDNode *N, const NVPTXSubtarget &Subtarget,
                           unsigned CodeAddrSpace, MachineFunction *F) {
   // We use ldg (i.e. ld.global.nc) for invariant loads from the global address
   // space.
   //
   // We have two ways of identifying invariant loads: Loads may be explicitly
   // marked as invariant, or we may infer them to be invariant.
   //
-  // We currently infer invariance only for kernel function pointer params that
-  // are noalias (i.e. __restrict) and never written to.
+  // We currently infer invariance for loads from
+  //  - constant global variables, and
+  //  - kernel function pointer params that are noalias (i.e. __restrict) and
+  //    never written to.
   //
   // TODO: Perform a more powerful invariance analysis (ideally IPO, and ideally
   // not during the SelectionDAG phase).
   //
   // TODO: Infer invariance only at -O2.  We still want to use ldg at -O0 for
   // explicitly invariant loads because these are how clang tells us to use ldg
   // when the user uses a builtin.
   if (!Subtarget.hasLDG() || CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL)
     return false;
 
   if (N->isInvariant())
     return true;
 
-  // Load wasn't explicitly invariant.  Attempt to infer invariance.
-  if (!isKernelFunction(F->getFunction()))
-    return false;
+  bool IsKernelFn = isKernelFunction(F->getFunction());
 
-  // We use GetUnderlyingObjects() here instead of
-  // GetUnderlyingObject() mainly because the former looks through phi
-  // nodes while the latter does not. We need to look through phi
-  // nodes to handle pointer induction variables.
+  // We use GetUnderlyingObjects() here instead of GetUnderlyingObject() mainly
+  // because the former looks through phi nodes while the latter does not. We
+  // need to look through phi nodes to handle pointer induction variables.
   SmallVector<Value *, 8> Objs;
   GetUnderlyingObjects(const_cast<Value *>(N->getMemOperand()->getValue()),
                        Objs, F->getDataLayout());
-  for (Value *Obj : Objs) {
-    auto *A = dyn_cast<const Argument>(Obj);
-    if (!A || !A->onlyReadsMemory() || !A->hasNoAliasAttr()) return false;
-  }
 
-  return true;
+  return all_of(Objs, [&](Value *V) {
+    if (auto *A = dyn_cast<const Argument>(V))
+      return IsKernelFn && A->onlyReadsMemory() && A->hasNoAliasAttr();
+    if (auto *GV = dyn_cast<const GlobalVariable>(V))
+      return GV->isConstant();
+    return false;
+  });
 }
 
 bool NVPTXDAGToDAGISel::tryIntrinsicNoChain(SDNode *N) {
   unsigned IID = cast<ConstantSDNode>(N->getOperand(0))->getZExtValue();
   switch (IID) {
   default:
     return false;
   case Intrinsic::nvvm_texsurf_handle_internal:
bool NVPTXDAGToDAGISel::tryLDGLDU(SDNode *N) {
   }
   InstVTs.push_back(MVT::Other);
   SDVTList InstVTList = CurDAG->getVTList(InstVTs);
 
   if (SelectDirectAddr(Op1, Addr)) {
     switch (N->getOpcode()) {
     default:
       return false;
+    case ISD::LOAD:
     case ISD::INTRINSIC_W_CHAIN:
       if (IsLDG)
         Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                      NVPTX::INT_PTX_LDG_GLOBAL_i8avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_i16avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_i32avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_i64avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_f16avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_f16x2avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_f32avar,
                                      NVPTX::INT_PTX_LDG_GLOBAL_f64avar);
       else
         Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                      NVPTX::INT_PTX_LDU_GLOBAL_i8avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_i16avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_i32avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_i64avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_f16avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_f16x2avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_f32avar,
                                      NVPTX::INT_PTX_LDU_GLOBAL_f64avar);
       break;
+    case NVPTXISD::LoadV2:
     case NVPTXISD::LDGV2:
       Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                    NVPTX::INT_PTX_LDG_G_v2i8_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2i16_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2i32_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2i64_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2f16_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2f16x2_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2f32_ELE_avar,
                                    NVPTX::INT_PTX_LDG_G_v2f64_ELE_avar);
       break;
     case NVPTXISD::LDUV2:
       Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                    NVPTX::INT_PTX_LDU_G_v2i8_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2i16_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2i32_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2i64_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2f16_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2f16x2_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2f32_ELE_avar,
                                    NVPTX::INT_PTX_LDU_G_v2f64_ELE_avar);
       break;
+    case NVPTXISD::LoadV4:
     case NVPTXISD::LDGV4:
       Opcode = pickOpcodeForVT(EltVT.getSimpleVT().SimpleTy,
                                NVPTX::INT_PTX_LDG_G_v4i8_ELE_avar,
                                NVPTX::INT_PTX_LDG_G_v4i16_ELE_avar,
                                NVPTX::INT_PTX_LDG_G_v4i32_ELE_avar, None,
                                NVPTX::INT_PTX_LDG_G_v4f16_ELE_avar,
                                NVPTX::INT_PTX_LDG_G_v4f16x2_ELE_avar,
                                NVPTX::INT_PTX_LDG_G_v4f32_ELE_avar, None);

llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll

+; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+
+; Check load from constant global variables.  These loads should be
+; ld.global.nc (aka ldg).
+
+@gv_float = external constant float
+@gv_float2 = external constant <2 x float>
+@gv_float4 = external constant <4 x float>
+
+; CHECK-LABEL: test_gv_float()
+define float @test_gv_float() {
+; CHECK: ld.global.nc.f32
+  %v = load float, float* @gv_float
+  ret float %v
+}
+
+; CHECK-LABEL: test_gv_float2()
+define <2 x float> @test_gv_float2() {
+  %v = load <2 x float>, <2 x float>* @gv_float2
+  ret <2 x float> %v
+}
+
+; CHECK-LABEL: test_gv_float4()
+define <4 x float> @test_gv_float4() {
+  %v = load <4 x float>, <4 x float>* @gv_float4
+  ret <4 x float> %v
+}