[PATCH] D106909: [clang] Add clang builtins support for gfx90a

[Matt Arsenault via Phabricator via cfe-commits]

arsenm added a comment.

Missing sema test for rejecting the builtins on targets that don't support this



================
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")
----------------
"_2f16" looks weird to me. The instruction names call it "pk"


================
Comment at: clang/lib/CodeGen/CGBuiltin.cpp:16114
+    Intrinsic::ID IID;
+    llvm::Type *ArgTy = llvm::Type::getDoubleTy(getLLVMContext());
+    switch (BuiltinID) {
----------------
Initializing this here is strange, sink down to the double case


================
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);
+}
----------------
If you're going to bother testing the ISA, is it worth testing rtn and no rtn versions?


Repository:
  rG LLVM Github Monorepo

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

https://reviews.llvm.org/D106909