arsenm added inline comments.
================
Comment at: clang/include/clang/Basic/BuiltinsAMDGPU.def:201
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_f32, "ff*1fi", "t",
"gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fadd_2f16, "hh*1hi", "t",
"gfx90a-insts")
+TARGET_BUILTIN(__builtin_amdgcn_global_atomic_fmin_f64, "dd*1di", "t",
"gfx90a-insts")
----------------
yaxunl wrote:
> arsenm wrote:
> > "_2f16" looks weird to me. The instruction names call it "pk"
> This is to have a consistent postfix naming convention, since the stem part
> here are the same. the postfix is for the argument type of the builtin
> function.
Just a plain 2 isn't consistent either. The llvm type naming convention would
add a v prefix, but the builtins should probably follow the instructions
================
Comment at: clang/lib/CodeGen/CGBuiltin.cpp:16114
+ Intrinsic::ID IID;
+ llvm::Type *ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
+ switch (BuiltinID) {
----------------
gandhi21299 wrote:
> arsenm wrote:
> > Initializing this here is strange, sink down to the double case
> You mean push it down after line 16119?
Yes
================
Comment at: clang/test/CodeGenOpenCL/builtins-fp-atomics.cl:112
+kernel void test_flat_global_max(__global double *addr, double x){
+ __builtin_amdgcn_flat_atomic_fmax_f64(addr, x, memory_order_relaxed);
+}
----------------
gandhi21299 wrote:
> arsenm wrote:
> > If you're going to bother testing the ISA, is it worth testing rtn and no
> > rtn versions?
> Sorry, what do you mean by rtn version?
Most atomics can be optimized if they don't return the in memory value if the
value is unused
Repository:
rG LLVM Github Monorepo
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D106909/new/
https://reviews.llvm.org/D106909
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
