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", 
+TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_fadd_v2bf16, "V2sV2s*3V2s", "t", 
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/
-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 ?

  rG LLVM Github Monorepo


cfe-commits mailing list

Reply via email to