[NVPTX] make load on global readonly memory to use ldg

[wengxt]

Summary

[NVPTX] make load on global readonly memory to use ldg

Summary:
As describe in [1], ld.global.nc may be used to load memory by nvcc when
restrict is used and compiler can detect whether read-only data cache
is safe to use.

This patch will try to check whether ldg is safe to use and use them to
replace ld.global when possible. This change can improve the performance
by 18~29% on affected kernels (ratt*_kernel and rwdot*_kernel) in
S3D benchmark of shoc [2].

Patched by Xuetian Weng.

[1] http://docs.nvidia.com/cuda/kepler-tuning-guide/#read-only-data-cache
[2] https://github.com/vetter/shoc

Test Plan: test/CodeGen/NVPTX/load-with-non-coherent-cache.ll

[jingyue]

"As describe in [1]"

I don't see the definition of [1].

You should mention the performance gain you got. SHOC is open-source.

[wengxt]

update commit message

[jingyue]

Can you mention the specific benchmarks (benchmark name or kernel name) and how much speedup they enjoy individually? This gives people something to refer to later. SHOC is open-source and frequently cited, so people would love to see in detail how your work affects that.

Otherwise, LGTM. Thanks!

lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

 //===-- NVPTXISelDAGToDAG.cpp - A dag to dag inst selector for NVPTX ------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines an instruction selector for the NVPTX target.
 //
 //===----------------------------------------------------------------------===//
 
 #include "NVPTXISelDAGToDAG.h"
+#include "llvm/Analysis/ValueTracking.h"
 #include "llvm/IR/GlobalValue.h"
 #include "llvm/IR/Instructions.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/Debug.h"
 #include "llvm/Support/ErrorHandling.h"
 #include "llvm/Support/raw_ostream.h"
 #include "llvm/Target/TargetIntrinsicInfo.h"
 
if (const PointerType *PT = dyn_cast<PointerType>(Src->getType())) {
     case llvm::ADDRESS_SPACE_PARAM: return NVPTX::PTXLdStInstCode::PARAM;
     case llvm::ADDRESS_SPACE_CONST: return NVPTX::PTXLdStInstCode::CONSTANT;
     default: break;
     }
   }
   return NVPTX::PTXLdStInstCode::GENERIC;
 }
 
+static bool canLowerToLDG(MemSDNode *N, const NVPTXSubtarget &Subtarget,
+                          unsigned codeAddrSpace, const DataLayout &DL) {
+  if (!Subtarget.hasLDG() || codeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL) {
+    return false;
+  }
+
+  // Check whether load operates on a readonly argument.
+  bool canUseLDG = false;
+  if (const Argument *A = dyn_cast<const Argument>(
+          GetUnderlyingObject(N->getMemOperand()->getValue(), DL)))
+    canUseLDG = A->onlyReadsMemory() && A->hasNoAliasAttr();
+
+  return canUseLDG;
+}
+
 SDNode *NVPTXDAGToDAGISel::SelectIntrinsicNoChain(SDNode *N) {
   unsigned IID = cast<ConstantSDNode>(N->getOperand(0))->getZExtValue();
   switch (IID) {
   default:
     return nullptr;
   case Intrinsic::nvvm_texsurf_handle_internal:
     return SelectTexSurfHandle(N);
   }
if (LD->isIndexed())
     return nullptr;
 
   if (!LoadedVT.isSimple())
     return nullptr;
 
   // Address Space Setting
   unsigned int codeAddrSpace = getCodeAddrSpace(LD);
 
+  if (canLowerToLDG(LD, *Subtarget, codeAddrSpace, CurDAG->getDataLayout())) {
+    return SelectLDGLDU(N);
+  }
+
   // Volatile Setting
   // - .volatile is only availalble for .global and .shared
   bool isVolatile = LD->isVolatile();
   if (codeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
       codeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
       codeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
     isVolatile = false;
 
SDNode *NVPTXDAGToDAGISel::SelectLoadVector(SDNode *N) {
   EVT LoadedVT = MemSD->getMemoryVT();
 
   if (!LoadedVT.isSimple())
     return nullptr;
 
   // Address Space Setting
   unsigned int CodeAddrSpace = getCodeAddrSpace(MemSD);
 
+  if (canLowerToLDG(MemSD, *Subtarget, CodeAddrSpace, CurDAG->getDataLayout())) {
+    return SelectLDGLDU(N);
+  }
+
   // Volatile Setting
   // - .volatile is only availalble for .global and .shared
   bool IsVolatile = MemSD->isVolatile();
   if (CodeAddrSpace != NVPTX::PTXLdStInstCode::GLOBAL &&
       CodeAddrSpace != NVPTX::PTXLdStInstCode::SHARED &&
       CodeAddrSpace != NVPTX::PTXLdStInstCode::GENERIC)
     IsVolatile = false;
 
if (SelectDirectAddr(Op1, Addr)) {
     SDValue Ops[] = { Addr, Chain };
     LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops);
   } else if (TM.is64Bit() ? SelectADDRri64(Op1.getNode(), Op1, Base, Offset)
                           : SelectADDRri(Op1.getNode(), Op1, Base, Offset)) {
     if (TM.is64Bit()) {
       switch (N->getOpcode()) {
       default:
         return nullptr;
+      case ISD::LOAD:
       case ISD::INTRINSIC_W_CHAIN:
         if (IsLDG) {
           switch (EltVT.getSimpleVT().SimpleTy) {
           default:
             return nullptr;
           case MVT::i8:
             Opcode = NVPTX::INT_PTX_LDG_GLOBAL_i8ari64;
             break;
if (TM.is64Bit()) {
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f32ari64;
             break;
           case MVT::f64:
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f64ari64;
             break;
           }
         }
         break;
+      case NVPTXISD::LoadV2:
       case NVPTXISD::LDGV2:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v2i8_ELE_ari64;
           break;
         case MVT::i16:
if (TM.is64Bit()) {
         case MVT::f32:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f32_ELE_ari64;
           break;
         case MVT::f64:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f64_ELE_ari64;
           break;
         }
         break;
+      case NVPTXISD::LoadV4:
       case NVPTXISD::LDGV4:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v4i8_ELE_ari64;
           break;
         case MVT::i16:
if (TM.is64Bit()) {
           break;
         }
         break;
       }
     } else {
       switch (N->getOpcode()) {
       default:
         return nullptr;
+      case ISD::LOAD:
       case ISD::INTRINSIC_W_CHAIN:
         if (IsLDG) {
           switch (EltVT.getSimpleVT().SimpleTy) {
           default:
             return nullptr;
           case MVT::i8:
             Opcode = NVPTX::INT_PTX_LDG_GLOBAL_i8ari;
             break;
if (TM.is64Bit()) {
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f32ari;
             break;
           case MVT::f64:
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f64ari;
             break;
           }
         }
         break;
+      case NVPTXISD::LoadV2:
       case NVPTXISD::LDGV2:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v2i8_ELE_ari32;
           break;
         case MVT::i16:
if (TM.is64Bit()) {
         case MVT::f32:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f32_ELE_ari32;
           break;
         case MVT::f64:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f64_ELE_ari32;
           break;
         }
         break;
+      case NVPTXISD::LoadV4:
       case NVPTXISD::LDGV4:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v4i8_ELE_ari32;
           break;
         case MVT::i16:
if (SelectDirectAddr(Op1, Addr)) {
     SDValue Ops[] = { Base, Offset, Chain };
 
     LD = CurDAG->getMachineNode(Opcode, DL, N->getVTList(), Ops);
   } else {
     if (TM.is64Bit()) {
       switch (N->getOpcode()) {
       default:
         return nullptr;
+      case ISD::LOAD:
       case ISD::INTRINSIC_W_CHAIN:
         if (IsLDG) {
           switch (EltVT.getSimpleVT().SimpleTy) {
           default:
             return nullptr;
           case MVT::i8:
             Opcode = NVPTX::INT_PTX_LDG_GLOBAL_i8areg64;
             break;
if (TM.is64Bit()) {
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f32areg64;
             break;
           case MVT::f64:
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f64areg64;
             break;
           }
         }
         break;
+      case NVPTXISD::LoadV2:
       case NVPTXISD::LDGV2:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v2i8_ELE_areg64;
           break;
         case MVT::i16:
if (TM.is64Bit()) {
         case MVT::f32:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f32_ELE_areg64;
           break;
         case MVT::f64:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f64_ELE_areg64;
           break;
         }
         break;
+      case NVPTXISD::LoadV4:
       case NVPTXISD::LDGV4:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v4i8_ELE_areg64;
           break;
         case MVT::i16:
if (TM.is64Bit()) {
           break;
         }
         break;
       }
     } else {
       switch (N->getOpcode()) {
       default:
         return nullptr;
+      case ISD::LOAD:
       case ISD::INTRINSIC_W_CHAIN:
         if (IsLDG) {
           switch (EltVT.getSimpleVT().SimpleTy) {
           default:
             return nullptr;
           case MVT::i8:
             Opcode = NVPTX::INT_PTX_LDG_GLOBAL_i8areg;
             break;
if (TM.is64Bit()) {
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f32areg;
             break;
           case MVT::f64:
             Opcode = NVPTX::INT_PTX_LDU_GLOBAL_f64areg;
             break;
           }
         }
         break;
+      case NVPTXISD::LoadV2:
       case NVPTXISD::LDGV2:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v2i8_ELE_areg32;
           break;
         case MVT::i16:
if (TM.is64Bit()) {
         case MVT::f32:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f32_ELE_areg32;
           break;
         case MVT::f64:
           Opcode = NVPTX::INT_PTX_LDU_G_v2f64_ELE_areg32;
           break;
         }
         break;
+      case NVPTXISD::LoadV4:
       case NVPTXISD::LDGV4:
         switch (EltVT.getSimpleVT().SimpleTy) {
         default:
           return nullptr;
         case MVT::i8:
           Opcode = NVPTX::INT_PTX_LDG_G_v4i8_ELE_areg32;
           break;
         case MVT::i16:

test/CodeGen/NVPTX/load-with-non-coherent-cache.ll

+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 | FileCheck -check-prefix=SM20 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck -check-prefix=SM35 %s
+
+target datalayout = "e-p:64:64:64-i1:8:8-i8:8:8-i16:16:16-i32:32:32-i64:64:64-f32:32:32-f64:64:64-v16:16:16-v32:32:32-v64:64:64-v128:128:128-n16:32:64"
+target triple = "nvptx64-unknown-unknown"
+
+; SM20-LABEL: .visible .entry foo1(
+; SM20: ld.global.f32
+; SM35-LABEL: .visible .entry foo1(
+; SM35: ld.global.nc.f32
+define void @foo1(float * noalias readonly %from, float * %to) {
+  %1 = load float, float * %from
+  store float %1, float * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo2(
+; SM20: ld.global.f64
+; SM35-LABEL: .visible .entry foo2(
+; SM35: ld.global.nc.f64
+define void @foo2(double * noalias readonly %from, double * %to) {
+  %1 = load double, double * %from
+  store double %1, double * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo3(
+; SM20: ld.global.u16
+; SM35-LABEL: .visible .entry foo3(
+; SM35: ld.global.nc.u16
+define void @foo3(i16 * noalias readonly %from, i16 * %to) {
+  %1 = load i16, i16 * %from
+  store i16 %1, i16 * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo4(
+; SM20: ld.global.u32
+; SM35-LABEL: .visible .entry foo4(
+; SM35: ld.global.nc.u32
+define void @foo4(i32 * noalias readonly %from, i32 * %to) {
+  %1 = load i32, i32 * %from
+  store i32 %1, i32 * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo5(
+; SM20: ld.global.u64
+; SM35-LABEL: .visible .entry foo5(
+; SM35: ld.global.nc.u64
+define void @foo5(i64 * noalias readonly %from, i64 * %to) {
+  %1 = load i64, i64 * %from
+  store i64 %1, i64 * %to
+  ret void
+}
+
+; i128 is non standard integer in nvptx64
+; SM20-LABEL: .visible .entry foo6(
+; SM20: ld.global.u64
+; SM20: ld.global.u64
+; SM35-LABEL: .visible .entry foo6(
+; SM35: ld.global.nc.u64
+; SM35: ld.global.nc.u64
+define void @foo6(i128 * noalias readonly %from, i128 * %to) {
+  %1 = load i128, i128 * %from
+  store i128 %1, i128 * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo7(
+; SM20: ld.global.v2.u8
+; SM35-LABEL: .visible .entry foo7(
+; SM35: ld.global.nc.v2.u8
+define void @foo7(<2 x i8> * noalias readonly %from, <2 x i8> * %to) {
+  %1 = load <2 x i8>, <2 x i8> * %from
+  store <2 x i8> %1, <2 x i8> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo8(
+; SM20: ld.global.v2.u16
+; SM35-LABEL: .visible .entry foo8(
+; SM35: ld.global.nc.v2.u16
+define void @foo8(<2 x i16> * noalias readonly %from, <2 x i16> * %to) {
+  %1 = load <2 x i16>, <2 x i16> * %from
+  store <2 x i16> %1, <2 x i16> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo9(
+; SM20: ld.global.v2.u32
+; SM35-LABEL: .visible .entry foo9(
+; SM35: ld.global.nc.v2.u32
+define void @foo9(<2 x i32> * noalias readonly %from, <2 x i32> * %to) {
+  %1 = load <2 x i32>, <2 x i32> * %from
+  store <2 x i32> %1, <2 x i32> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo10(
+; SM20: ld.global.v2.u64
+; SM35-LABEL: .visible .entry foo10(
+; SM35: ld.global.nc.v2.u64
+define void @foo10(<2 x i64> * noalias readonly %from, <2 x i64> * %to) {
+  %1 = load <2 x i64>, <2 x i64> * %from
+  store <2 x i64> %1, <2 x i64> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo11(
+; SM20: ld.global.v2.f32
+; SM35-LABEL: .visible .entry foo11(
+; SM35: ld.global.nc.v2.f32
+define void @foo11(<2 x float> * noalias readonly %from, <2 x float> * %to) {
+  %1 = load <2 x float>, <2 x float> * %from
+  store <2 x float> %1, <2 x float> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo12(
+; SM20: ld.global.v2.f64
+; SM35-LABEL: .visible .entry foo12(
+; SM35: ld.global.nc.v2.f64
+define void @foo12(<2 x double> * noalias readonly %from, <2 x double> * %to) {
+  %1 = load <2 x double>, <2 x double> * %from
+  store <2 x double> %1, <2 x double> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo13(
+; SM20: ld.global.v4.u8
+; SM35-LABEL: .visible .entry foo13(
+; SM35: ld.global.nc.v4.u8
+define void @foo13(<4 x i8> * noalias readonly %from, <4 x i8> * %to) {
+  %1 = load <4 x i8>, <4 x i8> * %from
+  store <4 x i8> %1, <4 x i8> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo14(
+; SM20: ld.global.v4.u16
+; SM35-LABEL: .visible .entry foo14(
+; SM35: ld.global.nc.v4.u16
+define void @foo14(<4 x i16> * noalias readonly %from, <4 x i16> * %to) {
+  %1 = load <4 x i16>, <4 x i16> * %from
+  store <4 x i16> %1, <4 x i16> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo15(
+; SM20: ld.global.v4.u32
+; SM35-LABEL: .visible .entry foo15(
+; SM35: ld.global.nc.v4.u32
+define void @foo15(<4 x i32> * noalias readonly %from, <4 x i32> * %to) {
+  %1 = load <4 x i32>, <4 x i32> * %from
+  store <4 x i32> %1, <4 x i32> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo16(
+; SM20: ld.global.v4.f32
+; SM35-LABEL: .visible .entry foo16(
+; SM35: ld.global.nc.v4.f32
+define void @foo16(<4 x float> * noalias readonly %from, <4 x float> * %to) {
+  %1 = load <4 x float>, <4 x float> * %from
+  store <4 x float> %1, <4 x float> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo17(
+; SM20: ld.global.v2.f64
+; SM20: ld.global.v2.f64
+; SM35-LABEL: .visible .entry foo17(
+; SM35: ld.global.nc.v2.f64
+; SM35: ld.global.nc.v2.f64
+define void @foo17(<4 x double> * noalias readonly %from, <4 x double> * %to) {
+  %1 = load <4 x double>, <4 x double> * %from
+  store <4 x double> %1, <4 x double> * %to
+  ret void
+}
+
+; SM20-LABEL: .visible .entry foo18(
+; SM20: ld.global.u64
+; SM35-LABEL: .visible .entry foo18(
+; SM35: ld.global.nc.u64
+define void @foo18(float ** noalias readonly %from, float ** %to) {
+  %1 = load float *, float ** %from
+  store float * %1, float ** %to
+  ret void
+}
+
+!nvvm.annotations = !{!1 ,!2 ,!3 ,!4 ,!5 ,!6, !7 ,!8 ,!9 ,!10 ,!11 ,!12, !13, !14, !15, !16, !17, !18}
+!1 = !{void (float *, float *)* @foo1, !"kernel", i32 1}
+!2 = !{void (double *, double *)* @foo2, !"kernel", i32 1}
+!3 = !{void (i16 *, i16 *)* @foo3, !"kernel", i32 1}
+!4 = !{void (i32 *, i32 *)* @foo4, !"kernel", i32 1}
+!5 = !{void (i64 *, i64 *)* @foo5, !"kernel", i32 1}
+!6 = !{void (i128 *, i128 *)* @foo6, !"kernel", i32 1}
+!7 = !{void (<2 x i8> *, <2 x i8> *)* @foo7, !"kernel", i32 1}
+!8 = !{void (<2 x i16> *, <2 x i16> *)* @foo8, !"kernel", i32 1}
+!9 = !{void (<2 x i32> *, <2 x i32> *)* @foo9, !"kernel", i32 1}
+!10 = !{void (<2 x i64> *, <2 x i64> *)* @foo10, !"kernel", i32 1}
+!11 = !{void (<2 x float> *, <2 x float> *)* @foo11, !"kernel", i32 1}
+!12 = !{void (<2 x double> *, <2 x double> *)* @foo12, !"kernel", i32 1}
+!13 = !{void (<4 x i8> *, <4 x i8> *)* @foo13, !"kernel", i32 1}
+!14 = !{void (<4 x i16> *, <4 x i16> *)* @foo14, !"kernel", i32 1}
+!15 = !{void (<4 x i32> *, <4 x i32> *)* @foo15, !"kernel", i32 1}
+!16 = !{void (<4 x float> *, <4 x float> *)* @foo16, !"kernel", i32 1}
+!17 = !{void (<4 x double> *, <4 x double> *)* @foo17, !"kernel", i32 1}
+!18 = !{void (float **, float **)* @foo18, !"kernel", i32 1}