[PATCH] D112466: [NVPTX] Drop memory references of LDG/LDU

Artem Belevich via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Thu Jan 12 10:41:10 PST 2023


tra accepted this revision.
tra added inline comments.
This revision is now accepted and ready to land.


================
Comment at: llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp:1674-1675
 
-  MachineMemOperand *MemRef = Mem->getMemOperand();
-  CurDAG->setNodeMemRefs(cast<MachineSDNode>(LD), {MemRef});
-
----------------
asavonic wrote:
> tra wrote:
> > The function also handles regular loads/stores (e.g. `NVPTXISD::LoadV4`) and those should still mark the node as a memory reference.
> Can you elaborate why we need to treat these nodes separately? A regular load can be lowered to LDG if it matches `canLowerToLDG`, but the final instruction is still LDG (aka `ld.global.nc`).
Sorry, my mistake. I only looked at the case constant and didn't pay attention that we're lowering it into an LDG instruction.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D112466



More information about the llvm-commits mailing list