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
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
