Changes

Summary

  1. [libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin (details)
  2. [ARM] Transforming memcpy to Tail predicated Loop (details)
  3. [mlir] Update dstNode after DenseMap insertion in loop fusion pass. (details)
  4. [libomptarget][nfc] Refactor amdgpu partial barrier to simplify adding a second one (details)
  5. [mlir][tosa] Added div op, variadic concat. Removed placeholder. Spec v0.22 alignment. (details)
  6. [AMDGPU] Expose __builtin_amdgcn_perm for v_perm_b32 (details)
  7. [llvm][TextAPI] add mapping from OS string to Platform (details)
  8. [dfsan] Rename and fix an internal test issue for mmap+calloc (details)
Commit 7e9351b9dee225b9ab12ee9bbfc9f8b96ddd1a1d by jonathanchesterfield
[libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin

[libomptarget][amdgpu][nfc] Remove dead code from amdgpu plugin

Drops an enum that was identical to a HSA one, localises some functions where
they were only called from one TU. Covers everything internalize + adce can
identify as dead, except for msgpack::dump which is useful when debugging.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D102014
The file was modifiedopenmp/libomptarget/plugins/amdgpu/impl/data.cpp
The file was modifiedopenmp/libomptarget/plugins/amdgpu/impl/atmi.h
The file was modifiedopenmp/libomptarget/plugins/amdgpu/impl/internal.h
The file was modifiedopenmp/libomptarget/plugins/amdgpu/impl/system.cpp
The file was modifiedopenmp/libomptarget/plugins/amdgpu/src/rtl.cpp
Commit 9ff38e2d9dd791383fbaa80e02d65e9c1f0463ff by malhar.jajoo
[ARM] Transforming memcpy to Tail predicated Loop

This patch converts llvm.memcpy intrinsic into Tail Predicated
Hardware loops for a target that supports the Arm M-profile
Vector Extension (MVE).

From an implementation point of view, the patch

- adds an ARM specific SDAG Node (to which the llvm.memcpy intrinsic is lowered to, during first phase of ISel)
- adds a corresponding TableGen entry to generate a pseudo instruction, with a custom inserter,
  on matching the above node.
- Adds a custom inserter function that expands the pseudo instruction into MIR suitable
   to be (by later passes) into a WLSTP loop.

Reviewed By: dmgreen

Differential Revision: https://reviews.llvm.org/D99723
The file was modifiedllvm/lib/Target/ARM/ARMISelLowering.h
The file was modifiedllvm/lib/Target/ARM/ARMISelLowering.cpp
The file was modifiedllvm/test/CodeGen/Thumb2/LowOverheadLoops/memcall.ll
The file was modifiedllvm/lib/Target/ARM/ARMSelectionDAGInfo.cpp
The file was modifiedllvm/lib/Target/ARM/ARMTargetTransformInfo.h
The file was addedllvm/test/CodeGen/Thumb2/mve-tp-loop.mir
The file was modifiedllvm/lib/Target/ARM/ARMSubtarget.h
The file was modifiedllvm/lib/Target/ARM/ARMInstrMVE.td
The file was addedllvm/test/CodeGen/Thumb2/mve-tp-loop.ll
Commit 5dc1ed3f627ecfad119ada84d3534cc21b80f810 by sergei.grechanik
[mlir] Update dstNode after DenseMap insertion in loop fusion pass.

Reviewed By: vinayaka-polymage

Differential Revision: https://reviews.llvm.org/D101794
The file was modifiedmlir/test/Transforms/loop-fusion.mlir
The file was modifiedmlir/lib/Transforms/LoopFusion.cpp
Commit 44ee974e2f3ef120e1890d8aafb02fedc3c135e9 by jonathanchesterfield
[libomptarget][nfc] Refactor amdgpu partial barrier to simplify adding a second one

[libomptarget][nfc] Refactor amdgpu partial barrier to simplify adding a second one

D101976 would require a second barrier instance. This NFC to amdgpu makes it
simpler to add one (an extra global, one more line in init). Also renames the
current barrier to L0.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D102016
The file was modifiedopenmp/libomptarget/deviceRTLs/amdgcn/src/target_impl.hip
Commit d3e987c38917758fe5cf1e27ed26d08f0e9c0fe3 by rob.suderman
[mlir][tosa] Added div op, variadic concat. Removed placeholder. Spec v0.22 alignment.

Nearly complete alignment to spec v0.22
- Adds Div op
- Concat inputs now variadic
- Removes Placeholder op

Note: TF side PR https://github.com/tensorflow/tensorflow/pull/48921 deletes Concat legalizations to avoid breaking TensorFlow CI. This must be merged only after the TF PR has merged.

Reviewed By: rsuderman

Differential Revision: https://reviews.llvm.org/D101958
The file was modifiedmlir/include/mlir/Dialect/Tosa/IR/TosaTypesBase.td
The file was modifiedmlir/include/mlir/Dialect/Tosa/IR/TosaOps.td
The file was modifiedmlir/test/Dialect/Tosa/ops.mlir
Commit c714d037857f9c8e3bbe32e22ec22279121c378d by Stanislav.Mekhanoshin
[AMDGPU] Expose __builtin_amdgcn_perm for v_perm_b32

Differential Revision: https://reviews.llvm.org/D102022
The file was addedllvm/test/CodeGen/AMDGPU/llvm.amdgcn.perm.ll
The file was modifiedllvm/include/llvm/IR/IntrinsicsAMDGPU.td
The file was modifiedllvm/lib/Target/AMDGPU/AMDGPUInstrInfo.td
The file was modifiedclang/test/SemaOpenCL/builtins-amdgcn-error-vi.cl
The file was modifiedllvm/lib/Target/AMDGPU/SIISelLowering.cpp
The file was modifiedclang/test/CodeGenOpenCL/builtins-amdgcn-vi.cl
The file was modifiedclang/include/clang/Basic/BuiltinsAMDGPU.def
The file was modifiedllvm/lib/Target/AMDGPU/AMDGPURegisterBankInfo.cpp
Commit c4ed142e695f14ba5675ec6d12226ee706329a0f by Cyndy Ishida
[llvm][TextAPI] add mapping from OS string to Platform

* add utility for matching target triple OS value strings  to PlatformKind

This was reviewed offline by ributzka, steven_wu
The file was modifiedllvm/lib/TextAPI/Platform.cpp
The file was modifiedllvm/unittests/TextAPI/TextStubV4Tests.cpp
The file was modifiedllvm/include/llvm/TextAPI/Platform.h
Commit 87a6325fbe4315f3f24555797f216e96539a9397 by jianzhouzh
[dfsan] Rename and fix an internal test issue for mmap+calloc

The linker suggests using -Wl,-z,notext.

Replaced assert by exit also fixed this.

After renaming, interceptor.c would be used to test interceptors in general by D101204.

Reviewed By: morehouse

Differential Revision: https://reviews.llvm.org/D101649
The file was addedcompiler-rt/test/dfsan/mmap_at_init.c
The file was removedcompiler-rt/test/dfsan/interceptors.c