[PATCH] D146701: [AMDGPU] Create Subtarget Features for some of 16 bits atomic fadd instructions

Thu Mar 23 04:47:26 PDT 2023

mariusz-sikora-at-amd added inline comments.

================
Comment at: clang/include/clang/Basic/BuiltinsAMDGPU.def:233
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_v2bf16, "V2sV2s*1V2s", "t", "atomic-global-pk-add-bf16-inst")
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", "atomic-ds-pk-add-16-insts")

----------------
foad wrote:
> So __builtin_amdgcn_ds_atomic_fadd_v2f16 is missing here? (Just curious- I am not asking you to add it in this patch.)
I have a different change which will add only __builtin_amdgcn_ds_atomic_fadd_v2f16 (not pushed yet, because it depends on this change).

================
Comment at: llvm/lib/Target/AMDGPU/BUFInstructions.td:2889

-defm BUFFER_ATOMIC_ADD_F32    : MUBUF_Real_Atomic_vi <0x4d>;
+let SubtargetPredicate = HasAtomicBufferGlobalPkAddF16NoRtnInsts in {
 defm BUFFER_ATOMIC_PK_ADD_F16 : MUBUF_Real_Atomic_vi <0x4e>;
----------------
foad wrote:
> Could remove the braces if you prefer - then you don't need the "End" comment either.
So, as I understand from other comment:

> Generally Real instructions copy their predicates from the corresponding Pseudo, so this should not be required here. Please check the other places where you have added predicates to Real instructions too.

We do not need this (L2889) Predicate, because it was added to Pseudo Instruction ?

Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D146701