[PATCH] D100394: [Clang][NVPTX] Add NVPTX intrinsics and builtins for CUDA PTX cp.async instructions

Stuart Adams via Phabricator via llvm-commits llvm-commits at lists.llvm.org
Tue Apr 20 06:38:23 PDT 2021


nyalloc updated this revision to Diff 338848.
nyalloc added a comment.

Addressed @tra's review comments.

- `_b64` postfix is removed from `mbarrier` intrinsics and builtins.
- `CHECK-LABEL` is introduced in `builtins-nvptx.c`
- Code style updated in `IntrinsicsNVVM.td`: `llvm_globali8ptr_ty` -> `llvm_global_i8ptr_ty` etc
- `int_nvvm_mbarrier_inval` && `int_nvvm_mbarrier_inval_shared` marked as `IntrWriteMem, IntrArgMemOnly, WriteOnly<ArgIndex<0>>, NoCapture<ArgIndex<0>>`
- `int_nvvm_mbarrier_pending_count` marked as `IntrNoMem`
- `ALL` check prefix added to `async-copy.ll`, `ALL-LABEL` used appropriately


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

https://reviews.llvm.org/D100394

Files:
  clang/include/clang/Basic/BuiltinsNVPTX.def
  clang/test/CodeGen/builtins-nvptx.c
  llvm/include/llvm/IR/IntrinsicsNVVM.td
  llvm/lib/Target/NVPTX/NVPTXInstrInfo.td
  llvm/lib/Target/NVPTX/NVPTXIntrinsics.td
  llvm/test/CodeGen/NVPTX/async-copy.ll
  llvm/test/CodeGen/NVPTX/mbarrier.ll

-------------- next part --------------
A non-text attachment was scrubbed...
Name: D100394.338848.patch
Type: text/x-patch
Size: 35462 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/llvm-commits/attachments/20210420/a97ae856/attachment.bin>


More information about the llvm-commits mailing list