[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