[PATCH] D112466: [NVPTX] Annotate LDG/LDU instructions as mayLoad

Mon Oct 25 11:16:15 PDT 2021

asavonic updated this revision to Diff 382062.
asavonic added a comment.

That makes sense, thanks a lot for the explanation!
Machine verifier checks for memoperands to determine whether an instruction is a load or a store. Perhaps we can just drop them for LDG/LDU if we don't want them to be counted as loads?


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D112466/new/

https://reviews.llvm.org/D112466

Files:
  llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
  llvm/test/CodeGen/NVPTX/bug26185-2.ll
  llvm/test/CodeGen/NVPTX/bug26185.ll
  llvm/test/CodeGen/NVPTX/ldg-invariant.ll
  llvm/test/CodeGen/NVPTX/ldu-i8.ll
  llvm/test/CodeGen/NVPTX/ldu-ldg.ll
  llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll
  llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
  llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll


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

--- llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll
+++ llvm/test/CodeGen/NVPTX/read-global-variable-constant.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Check load from constant global variables.  These loads should be
 ; ld.global.nc (aka ldg).
Index: llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
+++ llvm/test/CodeGen/NVPTX/load-with-non-coherent-cache.ll
@@ -1,5 +1,5 @@
-; 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
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_20 -verify-machineinstrs | FileCheck -check-prefix=SM20 %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | 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"
Index: llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll
+++ llvm/test/CodeGen/NVPTX/ldu-reg-plus-offset.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
 
 target datalayout = "e-p:32:32:32-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"
 
Index: llvm/test/CodeGen/NVPTX/ldu-ldg.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/ldu-ldg.ll
+++ llvm/test/CodeGen/NVPTX/ldu-ldg.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
 
 
 declare i8 @llvm.nvvm.ldu.global.i.i8.p1i8(i8 addrspace(1)* %ptr, i32 %align)
Index: llvm/test/CodeGen/NVPTX/ldu-i8.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/ldu-i8.ll
+++ llvm/test/CodeGen/NVPTX/ldu-i8.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_20 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_20 -verify-machineinstrs | FileCheck %s
 
 target datalayout = "e-p:32:32:32-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"
 
Index: llvm/test/CodeGen/NVPTX/ldg-invariant.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/ldg-invariant.ll
+++ llvm/test/CodeGen/NVPTX/ldg-invariant.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx64 -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Check that invariant loads from the global addrspace are lowered to
 ; ld.global.nc.
Index: llvm/test/CodeGen/NVPTX/bug26185.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/bug26185.ll
+++ llvm/test/CodeGen/NVPTX/bug26185.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Verify that we correctly emit code for i8 ldg/ldu. We do not expose 8-bit
 ; registers in the backend, so these loads need special handling.
Index: llvm/test/CodeGen/NVPTX/bug26185-2.ll
===================================================================
--- llvm/test/CodeGen/NVPTX/bug26185-2.ll
+++ llvm/test/CodeGen/NVPTX/bug26185-2.ll
@@ -1,4 +1,4 @@
-; RUN: llc < %s -march=nvptx -mcpu=sm_35 | FileCheck %s
+; RUN: llc < %s -march=nvptx -mcpu=sm_35 -verify-machineinstrs | FileCheck %s
 
 ; Verify that we correctly emit code for extending ldg/ldu. We do not expose
 ; extending variants in the backend, but the ldg/ldu selection code may pick
Index: llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
===================================================================
--- llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
+++ llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp
@@ -1657,9 +1657,6 @@
     LD = CurDAG->getMachineNode(Opcode.getValue(), DL, InstVTList, Ops);
   }
 
-  MachineMemOperand *MemRef = Mem->getMemOperand();
-  CurDAG->setNodeMemRefs(cast<MachineSDNode>(LD), {MemRef});
-
   // For automatic generation of LDG (through SelectLoad[Vector], not the
   // intrinsics), we may have an extending load like:
   //


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D112466.382062.patch
Type: text/x-patch
Size: 4794 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20211025/4f7e0884/attachment.bin>
