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

[Jay Foad via Phabricator via llvm-commits]

foad 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")
 
----------------
So __builtin_amdgcn_ds_atomic_fadd_v2f16 is missing here? (Just curious- I am not asking you to add it in this patch.)


================
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>;
----------------
Could remove the braces if you prefer - then you don't need the "End" comment either.


================
Comment at: llvm/lib/Target/AMDGPU/FLATInstructions.td:1916
+let SubtargetPredicate = HasAtomicBufferGlobalPkAddF16NoRtnInsts in
+defm GLOBAL_ATOMIC_PK_ADD_F16 : FLAT_Global_Real_Atomics_vi <0x04e, 0>;
+
----------------
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.


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D146701