AMDGPU: Verify subtarget specific builtins

[arsenm]

Summary

Cleanup setup of subtarget features.

[echristo]

lib/Basic/Targets.cpp

1795

Extra whitespace.

2010–2014

This is typically more of the "move the cpu checks down here" area from what you'd have above. Also you're not calling the target independent version of initFeatureMap - is that done on purpose?

[arsenm]

lib/Basic/Targets.cpp

2010–2014

Not sure what you mean exactly by "move the cpu checks down here". Do you mean setting all of the feature has* member variables should be moved out of the AMDGPUTargetInfo constructor + setCPU into here?

That was not done on purpose. I'm pretty confused about what all of these functions are for. I kind of expected all of this to automatically work from the target's set of defined features known from the backend. There seem to be a set of functions for manually parsing user specified features on a function and for those implied by the subtarget. Why is initFeatureMap separate from handleTargetFeatures? initFeatureMap seems to be almost the same thing with the addition of the CPU features. It's not clear to me what the difference between hasFeature and validateCpuSupports is supposed to mean, or why setFeatureEnabled is virtual.

[echristo]

Replied inline, I hope this is helpful :)

lib/Basic/Targets.cpp

2010–2014

Attempting to answer your questions, I'm mostly going phrase by phrase :)

a) Yes, I mean this.
b) None of this code is connected to the backend at all. There's some disagreement as to what level the front end should connect to the backend in this way. It's probably better left for another day.
c) Yes, these functions are all part of handling features that can be set on a function, or are the defaults coming in via the driver.
d) They're separate, sadly, because I couldn't get them together in the time I spent rewriting it. Take a look at the comment at Targets.cpp:8086 for a bit more detail.
e) validateCpuSupports is only used in one case right now and that's for a particular builtin that's only supported on one target. Go ahead and ignore that for your purposes now, if you later want to support it we can talk about it.
f) I'm not sure why you care whether or not setFeatureEnabled is virtual, but we can discuss it if you'd like - and why you think it shouldn't be :)

Most of these functions have few (or one) use and it is in Targets.cpp. (Don't worry about the ones outside of Targets.cpp for now, that's part of the function attribute code which was the cause of the cleanup (I promise it was) and refactoring I did, but doesn't affect clang initialization.

[arsenm]

Try to move more code into initFeatureMap.

I'm not sure how the booleans for features in the class are for now. X86 seems to have them, but it seems they are only used with user specified features? The only ones that matter right now don't depend on the subtarget, so I've left that in the constructor

[arsenm]

ping

[echristo]

Seems pretty reasonable. If I missed anything we can pick it up later.

-eric

lib/Basic/Targets.cpp

2010

Means it was either passed in via -mcpu or is set on the function via an attribute.

[arsenm]

r272091

include/clang/Basic/BuiltinsAMDGPU.def

 //==- BuiltinsAMDGPU.def - AMDGPU Builtin function database ------*- C++ -*-==//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 //
 // This file defines the AMDGPU-specific builtin function database. Users of
 // this file must define the BUILTIN macro to make use of this information.
 //
 //===----------------------------------------------------------------------===//
 
 // The format of this database matches clang/Basic/Builtins.def.
 
+#if defined(BUILTIN) && !defined(TARGET_BUILTIN)
+#   define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE) BUILTIN(ID, TYPE, ATTRS)
+#endif
+
 BUILTIN(__builtin_amdgcn_s_barrier, "v", "n")
 BUILTIN(__builtin_amdgcn_div_scale, "dddbb*", "n")
 BUILTIN(__builtin_amdgcn_div_scalef, "fffbb*", "n")
 BUILTIN(__builtin_amdgcn_div_fmas, "ddddb", "nc")
 BUILTIN(__builtin_amdgcn_div_fmasf, "ffffb", "nc")
 BUILTIN(__builtin_amdgcn_div_fixup, "dddd", "nc")
 BUILTIN(__builtin_amdgcn_div_fixupf, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_trig_preop, "ddi", "nc")

 BUILTIN(__builtin_amdgcn_cubetc, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_cubema, "ffff", "nc")
 BUILTIN(__builtin_amdgcn_s_memtime, "LUi", "n")
 BUILTIN(__builtin_amdgcn_s_sleep, "vIi", "n")
 
 //===----------------------------------------------------------------------===//
 // VI+ only builtins.
 //===----------------------------------------------------------------------===//
-BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n")
+
+TARGET_BUILTIN(__builtin_amdgcn_s_memrealtime, "LUi", "n", "s-memrealtime")
 
 //===----------------------------------------------------------------------===//
 // Legacy names with amdgpu prefix
 //===----------------------------------------------------------------------===//
 
 BUILTIN(__builtin_amdgpu_rsq, "dd", "nc")
 BUILTIN(__builtin_amdgpu_rsqf, "ff", "nc")
 BUILTIN(__builtin_amdgpu_ldexp, "ddi", "nc")
 BUILTIN(__builtin_amdgpu_ldexpf, "ffi", "nc")
 
 #undef BUILTIN
+#undef TARGET_BUILTIN

lib/Basic/Targets.cpp

static const char *const DataLayoutStringR600 =
   "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
   "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
 
 static const char *const DataLayoutStringSI =
   "e-p:32:32-p1:64:64-p2:64:64-p3:32:32-p4:64:64-p5:32:32-p24:64:64"
   "-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128"
   "-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64";
 
-class AMDGPUTargetInfo : public TargetInfo {
+class AMDGPUTargetInfo final : public TargetInfo {
   static const Builtin::Info BuiltinInfo[];
   static const char * const GCCRegNames[];
 
   /// \brief The GPU profiles supported by the AMDGPU target.
   enum GPUKind {
     GK_NONE,
     GK_R600,
     GK_R600_DOUBLE_OPS,
     GK_R700,
     GK_R700_DOUBLE_OPS,
     GK_EVERGREEN,
     GK_EVERGREEN_DOUBLE_OPS,
     GK_NORTHERN_ISLANDS,
     GK_CAYMAN,
     GK_SOUTHERN_ISLANDS,
     GK_SEA_ISLANDS,
     GK_VOLCANIC_ISLANDS
   } GPU;
 
   bool hasFP64:1;
   bool hasFMAF:1;
   bool hasLDEXPF:1;
 
+  static bool isAMDGCN(const llvm::Triple &TT) {
+    return TT.getArch() == llvm::Triple::amdgcn;
+  }
+
 public:
   AMDGPUTargetInfo(const llvm::Triple &Triple)
-    : TargetInfo(Triple) {
-
-    if (Triple.getArch() == llvm::Triple::amdgcn) {
-      resetDataLayout(DataLayoutStringSI);
-      GPU = GK_SOUTHERN_ISLANDS;
+    : TargetInfo(Triple) ,
+      GPU(isAMDGCN(Triple) ? GK_SOUTHERN_ISLANDS : GK_R600),
+      hasFP64(false),
+      hasFMAF(false),
+      hasLDEXPF(false) {
+    if (getTriple().getArch() == llvm::Triple::amdgcn) {
       hasFP64 = true;
       hasFMAF = true;
       hasLDEXPF = true;
-    } else {
-      resetDataLayout(DataLayoutStringR600);
-      GPU = GK_R600;
-      hasFP64 = false;
-      hasFMAF = false;
-      hasLDEXPF = false;
     }
+
+    resetDataLayout(getTriple().getArch() == llvm::Triple::amdgcn ?
+                    DataLayoutStringSI : DataLayoutStringR600);
+
     AddrSpaceMap = &AMDGPUAddrSpaceMap;
     UseAddrSpaceMapMangling = true;
   }
 
   uint64_t getPointerWidthV(unsigned AddrSpace) const override {
     if (GPU <= GK_CAYMAN)
       return 32;
 
bool validateAsmConstraint(const char *&Name,
     case 'v': // vgpr
     case 's': // sgpr
       Info.setAllowsRegister();
       return true;
     }
     return false;
   }
 
+  bool initFeatureMap(llvm::StringMap<bool> &Features,
+                      DiagnosticsEngine &Diags, StringRef CPU,
+                      const std::vector<std::string> &FeatureVec) const override;
+
   ArrayRef<Builtin::Info> getTargetBuiltins() const override {
     return llvm::makeArrayRef(BuiltinInfo,
                         clang::AMDGPU::LastTSBuiltin - Builtin::FirstTSBuiltin);
   }
 
   bool isCLZForZeroUndef() const override {
     // It is -1 instead of expected for intrinsic.
     return true;
if (Opts.OpenCL) {
       }
     }
   }
 
   BuiltinVaListKind getBuiltinVaListKind() const override {
     return TargetInfo::CharPtrBuiltinVaList;
   }
 
-  bool setCPU(const std::string &Name) override {
-    GPU = llvm::StringSwitch<GPUKind>(Name)
+  static GPUKind parseR600Name(StringRef Name) {
+    return llvm::StringSwitch<GPUKind>(Name)
       .Case("r600" ,    GK_R600)
       .Case("rv610",    GK_R600)
       .Case("rv620",    GK_R600)
       .Case("rv630",    GK_R600)
       .Case("rv635",    GK_R600)
       .Case("rs780",    GK_R600)
       .Case("rs880",    GK_R600)
       .Case("rv670",    GK_R600_DOUBLE_OPS)
return llvm::StringSwitch<GPUKind>(Name)
       .Case("juniper",  GK_EVERGREEN)
       .Case("hemlock",  GK_EVERGREEN_DOUBLE_OPS)
       .Case("cypress",  GK_EVERGREEN_DOUBLE_OPS)
       .Case("barts",    GK_NORTHERN_ISLANDS)
       .Case("turks",    GK_NORTHERN_ISLANDS)
       .Case("caicos",   GK_NORTHERN_ISLANDS)
       .Case("cayman",   GK_CAYMAN)
       .Case("aruba",    GK_CAYMAN)
+      .Default(GK_NONE);
+  }
+
+  static GPUKind parseAMDGCNName(StringRef Name) {
+    return llvm::StringSwitch<GPUKind>(Name)
       .Case("tahiti",   GK_SOUTHERN_ISLANDS)
       .Case("pitcairn", GK_SOUTHERN_ISLANDS)
       .Case("verde",    GK_SOUTHERN_ISLANDS)
       .Case("oland",    GK_SOUTHERN_ISLANDS)
       .Case("hainan",   GK_SOUTHERN_ISLANDS)
       .Case("bonaire",  GK_SEA_ISLANDS)
       .Case("kabini",   GK_SEA_ISLANDS)
       .Case("kaveri",   GK_SEA_ISLANDS)
       .Case("hawaii",   GK_SEA_ISLANDS)
       .Case("mullins",  GK_SEA_ISLANDS)
       .Case("tonga",    GK_VOLCANIC_ISLANDS)
       .Case("iceland",  GK_VOLCANIC_ISLANDS)
       .Case("carrizo",  GK_VOLCANIC_ISLANDS)
       .Case("fiji",     GK_VOLCANIC_ISLANDS)
       .Case("stoney",   GK_VOLCANIC_ISLANDS)
       .Default(GK_NONE);
-
-    if (GPU == GK_NONE) {
-      return false;
-    }
 
-    // Set the correct data layout
-    switch (GPU) {
-    case GK_NONE:
-    case GK_R600:
-    case GK_R700:
-    case GK_EVERGREEN:
-    case GK_NORTHERN_ISLANDS:
-      resetDataLayout(DataLayoutStringR600);
-      hasFP64 = false;
-      hasFMAF = false;
-      hasLDEXPF = false;
-      break;
-    case GK_R600_DOUBLE_OPS:
-    case GK_R700_DOUBLE_OPS:
-    case GK_EVERGREEN_DOUBLE_OPS:
-    case GK_CAYMAN:
-      resetDataLayout(DataLayoutStringR600);
-      hasFP64 = true;
-      hasFMAF = true;
-      hasLDEXPF = false;
-      break;
-    case GK_SOUTHERN_ISLANDS:
-    case GK_SEA_ISLANDS:
-    case GK_VOLCANIC_ISLANDS:
-      resetDataLayout(DataLayoutStringSI);
-      hasFP64 = true;
-      hasFMAF = true;
-      hasLDEXPF = true;
-      break;
-    }
+  bool setCPU(const std::string &Name) override {
+    if (getTriple().getArch() == llvm::Triple::amdgcn)
+      GPU = parseAMDGCNName(Name);
+    else
+      GPU = parseR600Name(Name);
 
-    return true;
+    return GPU != GK_NONE;
   }
 };
 
 const Builtin::Info AMDGPUTargetInfo::BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
+#define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE)                               \
+  { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, FEATURE },
 #include "clang/Basic/BuiltinsAMDGPU.def"
 };
 const char * const AMDGPUTargetInfo::GCCRegNames[] = {
   "v0", "v1", "v2", "v3", "v4", "v5", "v6", "v7",
   "v8", "v9", "v10", "v11", "v12", "v13", "v14", "v15",
   "v16", "v17", "v18", "v19", "v20", "v21", "v22", "v23",
   "v24", "v25", "v26", "v27", "v28", "v29", "v30", "v31",
   "v32", "v33", "v34", "v35", "v36", "v37", "v38", "v39",
const char * const AMDGPUTargetInfo::GCCRegNames[] = {
   "exec", "vcc", "scc", "m0", "flat_scratch", "exec_lo", "exec_hi",
   "vcc_lo", "vcc_hi", "flat_scratch_lo", "flat_scratch_hi"
 };
 
 ArrayRef<const char *> AMDGPUTargetInfo::getGCCRegNames() const {
   return llvm::makeArrayRef(GCCRegNames);
 }
 
+bool AMDGPUTargetInfo::initFeatureMap(
+  llvm::StringMap<bool> &Features,
+  DiagnosticsEngine &Diags, StringRef CPU,
+  const std::vector<std::string> &FeatureVec) const {
+
+  // XXX - What does the member GPU mean if device name string passed here?
+  if (getTriple().getArch() == llvm::Triple::amdgcn) {
+    if (CPU.empty())
+      CPU = "tahiti";
+
+    switch (parseAMDGCNName(CPU)) {
+    case GK_SOUTHERN_ISLANDS:
+    case GK_SEA_ISLANDS:
+      break;
+
+    case GK_VOLCANIC_ISLANDS:
+      Features["s-memrealtime"] = true;
+      Features["16-bit-insts"] = true;
+      break;
+
+    case GK_NONE:
+      return false;
+    default:
+      llvm_unreachable("unhandled subtarget");
+    }
+  } else {
+    if (CPU.empty())
+      CPU = "r600";
+
+    switch (parseR600Name(CPU)) {
+    case GK_R600:
+    case GK_R700:
+    case GK_EVERGREEN:
+    case GK_NORTHERN_ISLANDS:
+      break;
+    case GK_R600_DOUBLE_OPS:
+    case GK_R700_DOUBLE_OPS:
+    case GK_EVERGREEN_DOUBLE_OPS:
+    case GK_CAYMAN:
+      Features["fp64"] = true;
+      break;
+    case GK_NONE:
+      return false;
+    default:
+      llvm_unreachable("unhandled subtarget");
+    }
+  }
+
+  return TargetInfo::initFeatureMap(Features, Diags, CPU, FeatureVec);
+}
+
 // Namespace for x86 abstract base class
 const Builtin::Info BuiltinInfo[] = {
 #define BUILTIN(ID, TYPE, ATTRS)                                               \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, nullptr },
 #define LIBBUILTIN(ID, TYPE, ATTRS, HEADER)                                    \
   { #ID, TYPE, ATTRS, HEADER, ALL_LANGUAGES, nullptr },
 #define TARGET_BUILTIN(ID, TYPE, ATTRS, FEATURE)                               \
   { #ID, TYPE, ATTRS, nullptr, ALL_LANGUAGES, FEATURE },

lib/Driver/Tools.cpp

static std::string getCPUName(const ArgList &Args, const llvm::Triple &T,
 
   case llvm::Triple::hexagon:
     return "hexagon" +
            toolchains::HexagonToolChain::GetTargetCPUVersion(Args).str();
 
   case llvm::Triple::systemz:
     return getSystemZTargetCPU(Args);
 
-  case llvm::Triple::r600:
   case llvm::Triple::amdgcn:
+    return getAMDGCNTargetGPU(Args);
+  case llvm::Triple::r600:
     return getR600TargetGPU(Args);
 
   case llvm::Triple::wasm32:
   case llvm::Triple::wasm64:
     return getWebAssemblyTargetCPU(Args);
   }
 }
 

test/CodeGenOpenCL/builtins-amdgcn-error.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -target-cpu tahiti -verify -S -o - %s
+
+// FIXME: We only get one error if the functions are the other order in the
+// file.
+
+typedef unsigned long ulong;
+
+ulong test_s_memrealtime()
+{
+  return __builtin_amdgcn_s_memrealtime(); // expected-error {{'__builtin_amdgcn_s_memrealtime' needs target feature s-memrealtime}}
+}
+
+void test_s_sleep(int x)
+{
+  __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
+}
+

test/CodeGenOpenCL/builtins-amdgcn-vi.cl

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -S -emit-llvm -o - %s | FileCheck %s
+
+typedef unsigned long ulong;
+
+
+// CHECK-LABEL: @test_s_memrealtime
+// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
+void test_s_memrealtime(global ulong* out)
+{
+  *out = __builtin_amdgcn_s_memrealtime();
+}

test/CodeGenOpenCL/builtins-amdgcn.cl

 
 // CHECK-LABEL: @test_s_memtime
 // CHECK: call i64 @llvm.amdgcn.s.memtime()
 void test_s_memtime(global ulong* out)
 {
   *out = __builtin_amdgcn_s_memtime();
 }
 
-// CHECK-LABEL: @test_s_memrealtime
-// CHECK: call i64 @llvm.amdgcn.s.memrealtime()
-void test_s_memrealtime(global ulong* out)
-{
-  *out = __builtin_amdgcn_s_memrealtime();
-}
-
 // CHECK-LABEL: @test_s_sleep
 // CHECK: call void @llvm.amdgcn.s.sleep(i32 1)
 // CHECK: call void @llvm.amdgcn.s.sleep(i32 15)
 void test_s_sleep()
 {
   __builtin_amdgcn_s_sleep(1);
   __builtin_amdgcn_s_sleep(15);
 }

test/SemaOpenCL/builtins-amdgcn.cl

-// RUN: %clang_cc1 -triple amdgcn-unknown-amdhsa -fsyntax-only -verify %s
-
-void test_s_sleep(int x)
-{
-  __builtin_amdgcn_s_sleep(x); // expected-error {{argument to '__builtin_amdgcn_s_sleep' must be a constant integer}}
-}