Print out valid AMDGCN processor identifiers.

Author: AlexVlx

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -12445,7 +12445,7 @@ def warn_zero_as_null_pointer_constant : Warning<
   InGroup<DiagGroup<"zero-as-null-pointer-constant">>, DefaultIgnore;
 
 def warn_not_eliding_copy_on_return : Warning<
-  "not eliding copy on return">, 
+  "not eliding copy on return">,
   InGroup<DiagGroup<"nrvo">>, DefaultIgnore;
 
 def err_nullability_cs_multilevel : Error<
@@ -13347,6 +13347,8 @@ def err_amdgcn_processor_is_arg_not_literal
 def err_amdgcn_processor_is_arg_invalid_value
     : Error<"the argument to __builtin_amdgcn_processor_is must be a valid "
             "AMDGCN processor identifier; '%0' is not valid">;
+def note_amdgcn_processor_is_valid_options
+    : Note<"valid AMDGCN processor identifiers are: %0">;
 def err_amdgcn_is_invocable_arg_invalid_value
     : Error<"the argument to __builtin_amdgcn_is_invocable must be either a "
             "target agnostic builtin or an AMDGCN target specific builtin; `%0`"

clang/lib/Basic/Targets/SPIR.cpp

@@ -156,3 +156,8 @@ void SPIRV64AMDGCNTargetInfo::setAuxTarget(const TargetInfo *Aux) {
 bool SPIRV64AMDGCNTargetInfo::isValidCPUName(StringRef CPU) const {
   return AMDGPUTI.isValidCPUName(CPU);
 }
+
+void SPIRV64AMDGCNTargetInfo::fillValidCPUList(
+    SmallVectorImpl<StringRef> &Values) const {
+  return AMDGPUTI.fillValidCPUList(Values);
+}

clang/lib/Basic/Targets/SPIR.h

@@ -453,6 +453,7 @@ class LLVM_LIBRARY_VISIBILITY SPIRV64AMDGCNTargetInfo final
   // This is only needed for validating arguments passed to
   // __builtin_amdgcn_processor_is
   bool isValidCPUName(StringRef Name) const override;
+  void fillValidCPUList(SmallVectorImpl<StringRef> &Values) const override;
 };
 
 } // namespace targets

clang/lib/Sema/SemaAMDGPU.cpp

@@ -393,10 +393,18 @@ Expr *SemaAMDGPU::ExpandAMDGPUPredicateBI(CallExpr *CE) {
     }
 
     StringRef N = GFX->getString();
-    if (!Ctx.getTargetInfo().isValidCPUName(N) &&
-        (!Ctx.getAuxTargetInfo() ||
-         !Ctx.getAuxTargetInfo()->isValidCPUName(N))) {
+    const TargetInfo &TI = Ctx.getTargetInfo();
+    const TargetInfo *AuxTI = Ctx.getAuxTargetInfo();
+    if (!TI.isValidCPUName(N) && (!AuxTI || !AuxTI->isValidCPUName(N))) {
       Diag(Loc, diag::err_amdgcn_processor_is_arg_invalid_value) << N;
+      SmallVector<StringRef, 32> ValidList;
+      if (TI.getTriple().getVendor() == llvm::Triple::VendorType::AMD)
+        TI.fillValidCPUList(ValidList);
+      else if (AuxTI) // Since the BI is present it must be and AMDGPU triple.
+        AuxTI->fillValidCPUList(ValidList);
+      if (!ValidList.empty())
+        Diag(Loc, diag::note_amdgcn_processor_is_valid_options)
+            << llvm::join(ValidList, ", ");
       return nullptr;
     }
     if (Ctx.getTargetInfo().getTriple().isSPIRV()) {

clang/test/CodeGen/amdgpu-feature-builtins-invalid-use.cpp

@@ -28,6 +28,7 @@ void invalid_uses(int *p, int x, const __amdgpu_feature_predicate_t &lv,
 
 void invalid_invocations(int x, const char* str) {
     // CHECK: error: the argument to __builtin_amdgcn_processor_is must be a valid AMDGCN processor identifier; 'not_an_amdgcn_gfx_id' is not valid
+    // CHECK-DAG: note: valid AMDGCN processor identifiers are: {{.*}}gfx{{.*}}
     if (__builtin_amdgcn_processor_is("not_an_amdgcn_gfx_id")) return;
     // CHECK: error: the argument to __builtin_amdgcn_processor_is must be a string literal
     if (__builtin_amdgcn_processor_is(str)) return;