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

_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to