[SYCL] Add property-based device_has diagnostics

clang/include/clang/Basic/DiagnosticFrontendKinds.td

@@ -283,8 +283,8 @@ def err_avx_calling_convention : Error<warn_avx_calling_convention.Text>;
 
 def warn_sycl_device_has_aspect_mismatch
     : Warning<"function '%0' uses aspect '%1' not listed in its "
-              "'sycl::device_has' attribute">, BackendInfo,
-              InGroup<SyclAspectMismatch>;
+              "%select{'device_has' property|'sycl::device_has' attribute}2">,
+              BackendInfo, InGroup<SyclAspectMismatch>;
 def note_sycl_aspect_propagated_from_call
     : Note<"propagated from call to function '%0'">, BackendInfo;
 

clang/lib/CodeGen/CodeGenAction.cpp

@@ -867,7 +867,8 @@ void BackendConsumer::AspectMismatchDiagHandler(
   assert(LocCookie.isValid() &&
          "Invalid location for caller in aspect mismatch diagnostic");
   Diags.Report(LocCookie, diag::warn_sycl_device_has_aspect_mismatch)
-      << llvm::demangle(D.getFunctionName().str()) << D.getAspect();
+      << llvm::demangle(D.getFunctionName().str()) << D.getAspect()
+      << D.isFromDeviceHasAttribute();
   for (const std::pair<StringRef, unsigned> &CalleeInfo : D.getCallChain()) {
     LocCookie = SourceLocation::getFromRawEncoding(CalleeInfo.second);
     assert(LocCookie.isValid() &&

clang/lib/Sema/SemaDeclAttr.cpp

@@ -7819,7 +7819,8 @@ void Sema::CheckSYCLAddIRAttributesFunctionAttrConflicts(Decl *D) {
   for (const auto *Attr : std::vector<AttributeCommonInfo *>{
            D->getAttr<ReqdWorkGroupSizeAttr>(),
            D->getAttr<IntelReqdSubGroupSizeAttr>(),
-           D->getAttr<WorkGroupSizeHintAttr>()})
+           D->getAttr<WorkGroupSizeHintAttr>(),
+           D->getAttr<SYCLDeviceHasAttr>()})
     if (Attr)
       Diag(Attr->getLoc(), diag::warn_sycl_old_and_new_kernel_attributes)
           << Attr;

clang/test/SemaSYCL/attr-add-ir-attributes-function-conflict.cpp

@@ -1,8 +1,10 @@
-// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -verify %s
 
 // Tests that add_ir_attributes_function causes a warning when appearing with
 // potentially conflicting SYCL attributes.
 
+#include "sycl.hpp"
+
 constexpr const char AttrName1[] = "Attr1";
 constexpr const char AttrVal1[] = "Val1";
 
@@ -20,10 +22,13 @@ int main() {
   EmptyWrapper.kernel_single_task<class EK1>([]() [[sycl::reqd_work_group_size(1)]] {});
   EmptyWrapper.kernel_single_task<class EK2>([]() [[sycl::reqd_work_group_size(1,2)]] {});
   EmptyWrapper.kernel_single_task<class EK3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
-  EmptyWrapper.kernel_single_task<class EK1>([]() [[sycl::work_group_size_hint(1)]] {});
-  EmptyWrapper.kernel_single_task<class EK2>([]() [[sycl::work_group_size_hint(1,2)]] {});
-  EmptyWrapper.kernel_single_task<class EK3>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
+  EmptyWrapper.kernel_single_task<class EK4>([]() [[sycl::work_group_size_hint(1)]] {});
+  EmptyWrapper.kernel_single_task<class EK5>([]() [[sycl::work_group_size_hint(1,2)]] {});
+  EmptyWrapper.kernel_single_task<class EK6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
   EmptyWrapper.kernel_single_task<class EK7>([]() [[sycl::reqd_sub_group_size(1)]] {});
+  EmptyWrapper.kernel_single_task<class EK8>([]() [[sycl::device_has()]] {});
+  EmptyWrapper.kernel_single_task<class EK9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
+  EmptyWrapper.kernel_single_task<class EK10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 
   // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapper.kernel_single_task<class NEK1>([]() [[sycl::reqd_work_group_size(1)]] {});
@@ -32,11 +37,17 @@ int main() {
   // expected-warning@+1 {{kernel has both attribute 'reqd_work_group_size' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapper.kernel_single_task<class NEK3>([]() [[sycl::reqd_work_group_size(1,2,3)]] {});
   // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
-  NonemptyWrapper.kernel_single_task<class NEK1>([]() [[sycl::work_group_size_hint(1)]] {});
+  NonemptyWrapper.kernel_single_task<class NEK4>([]() [[sycl::work_group_size_hint(1)]] {});
   // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
-  NonemptyWrapper.kernel_single_task<class NEK2>([]() [[sycl::work_group_size_hint(1,2)]] {});
+  NonemptyWrapper.kernel_single_task<class NEK5>([]() [[sycl::work_group_size_hint(1,2)]] {});
   // expected-warning@+1 {{kernel has both attribute 'work_group_size_hint' and kernel properties; conflicting properties are ignored}}
-  NonemptyWrapper.kernel_single_task<class NEK3>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
+  NonemptyWrapper.kernel_single_task<class NEK6>([]() [[sycl::work_group_size_hint(1,2,3)]] {});
   // expected-warning@+1 {{kernel has both attribute 'reqd_sub_group_size' and kernel properties; conflicting properties are ignored}}
   NonemptyWrapper.kernel_single_task<class NEK7>([]() [[sycl::reqd_sub_group_size(1)]] {});
+  // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapper.kernel_single_task<class NEK8>([]() [[sycl::device_has()]] {});
+  // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapper.kernel_single_task<class NEK9>([]() [[sycl::device_has(sycl::aspect::cpu)]] {});
+  // expected-warning@+1 {{kernel has both attribute 'device_has' and kernel properties; conflicting properties are ignored}}
+  NonemptyWrapper.kernel_single_task<class NEK10>([]() [[sycl::device_has(sycl::aspect::cpu, sycl::aspect::gpu)]] {});
 }

llvm/include/llvm/IR/DiagnosticInfo.h

@@ -1120,30 +1120,32 @@ class DiagnosticInfoDontCall : public DiagnosticInfo {
 
 void diagnoseAspectsMismatch(const Function *F,
                              const SmallVector<Function *, 8> &CallChain,
-                             StringRef Aspect);
+                             StringRef Aspect, bool FromDeviceHasAttribute);
 
 // Diagnostic information for SYCL aspects usage mismatch.
 class DiagnosticInfoAspectsMismatch : public DiagnosticInfo {
   StringRef FunctionName;
   unsigned LocCookie;
   llvm::SmallVector<std::pair<StringRef, unsigned>, 8> CallChain;
   StringRef Aspect;
+  bool FromDeviceHasAttribute;
 
 public:
   DiagnosticInfoAspectsMismatch(
       StringRef FunctionName, unsigned LocCookie,
       const llvm::SmallVector<std::pair<StringRef, unsigned>, 8> &CallChain,
-      StringRef Aspect)
+      StringRef Aspect, bool FromDeviceHasAttribute)
       : DiagnosticInfo(DK_AspectMismatch, DiagnosticSeverity::DS_Warning),
         FunctionName(FunctionName), LocCookie(LocCookie), CallChain(CallChain),
-        Aspect(Aspect) {}
+        Aspect(Aspect), FromDeviceHasAttribute(FromDeviceHasAttribute) {}
   StringRef getFunctionName() const { return FunctionName; }
   unsigned getLocCookie() const { return LocCookie; }
   const llvm::SmallVector<std::pair<StringRef, unsigned>, 8> &
   getCallChain() const {
     return CallChain;
   }
   StringRef getAspect() const { return Aspect; }
+  bool isFromDeviceHasAttribute() const { return FromDeviceHasAttribute; }
   void print(DiagnosticPrinter &DP) const override;
   static bool classof(const DiagnosticInfo *DI) {
     return DI->getKind() == DK_AspectMismatch;

llvm/lib/IR/DiagnosticInfo.cpp

@@ -449,7 +449,8 @@ void DiagnosticInfoDontCall::print(DiagnosticPrinter &DP) const {
 
 void llvm::diagnoseAspectsMismatch(const Function *F,
                                    const SmallVector<Function *, 8> &CallChain,
-                                   StringRef Aspect) {
+                                   StringRef Aspect,
+                                   bool FromDeviceHasAttribute) {
   unsigned LocCookie = 0;
   if (MDNode *MD = F->getMetadata("srcloc"))
     LocCookie =
@@ -466,7 +467,7 @@ void llvm::diagnoseAspectsMismatch(const Function *F,
   }
 
   DiagnosticInfoAspectsMismatch D(F->getName(), LocCookie, LoweredCallChain,
-                                  Aspect);
+                                  Aspect, FromDeviceHasAttribute);
   F->getContext().diagnose(D);
 }
 

llvm/lib/SYCLLowerIR/SYCLPropagateAspectsUsage.cpp

@@ -246,29 +246,39 @@ AspectsSetTy getAspectsUsedByInstruction(const Instruction &I,
 using FunctionToAspectsMapTy = DenseMap<Function *, AspectsSetTy>;
 using CallGraphTy = DenseMap<Function *, SmallPtrSet<Function *, 8>>;
 
+// Finds the first function in a list that uses a given aspect. Returns nullptr
+// if none of the functions satisfy the criteria.
+Function *findFirstAspectUsageCallee(
+    const SmallPtrSetImpl<Function *> &Callees,
+    const FunctionToAspectsMapTy &AspectsMap, int Aspect,
+    SmallPtrSetImpl<const Function *> *Visited = nullptr) {
+  for (Function *Callee : Callees) {
+    if (Visited && !Visited->insert(Callee).second)
+      continue;
+
+    auto AspectIt = AspectsMap.find(Callee);
+    if (AspectIt != AspectsMap.end() && AspectIt->second.contains(Aspect))
+      return Callee;
+  }
+  return nullptr;
+}
+
 // Constructs an aspect usage chain for a given aspect from the function to the
 // last callee in the first found chain.
 void constructAspectUsageChain(const Function *F,
                                const FunctionToAspectsMapTy &AspectsMap,
                                const CallGraphTy &CG, int Aspect,
-                               SmallVector<Function *, 8> &CallChain,
-                               SmallPtrSet<const Function *, 16> &Visited) {
+                               SmallVectorImpl<Function *> &CallChain,
+                               SmallPtrSetImpl<const Function *> &Visited) {
   const auto EdgeIt = CG.find(F);
   if (EdgeIt == CG.end())
     return;
 
-  for (Function *Callee : EdgeIt->second) {
-    if (!Visited.insert(Callee).second)
-      continue;
-
-    auto AspectIt = AspectsMap.find(Callee);
-    if (AspectIt == AspectsMap.end() || !AspectIt->second.contains(Aspect))
-      continue;
-
-    CallChain.push_back(Callee);
-    constructAspectUsageChain(Callee, AspectsMap, CG, Aspect, CallChain,
-                              Visited);
-    break;
+  if (Function *AspectUsingCallee = findFirstAspectUsageCallee(
+          EdgeIt->second, AspectsMap, Aspect, &Visited)) {
+    CallChain.push_back(AspectUsingCallee);
+    constructAspectUsageChain(AspectUsingCallee, AspectsMap, CG, Aspect,
+                              CallChain, Visited);
   }
 }
 
@@ -313,22 +323,33 @@ void validateUsedAspectsForFunctions(const FunctionToAspectsMapTy &Map,
       continue;
 
     Function *F = It.first;
-
-    // Entry points will have their declared aspects from their kernel call.
-    // To avoid double warnings, we skip them.
-    if (std::find(EntryPoints.begin(), EntryPoints.end(), F) !=
-        EntryPoints.end())
-      continue;
-
-    const MDNode *DeviceHasMD = F->getMetadata("sycl_declared_aspects");
-    if (!DeviceHasMD)
-      continue;
-
     AspectsSetTy DeviceHasAspectSet;
-    for (size_t I = 0; I != DeviceHasMD->getNumOperands(); ++I) {
-      const auto *CAM = cast<ConstantAsMetadata>(DeviceHasMD->getOperand(I));
-      const Constant *C = CAM->getValue();
-      DeviceHasAspectSet.insert(cast<ConstantInt>(C)->getSExtValue());
+    bool OriginatedFromAttribute = true;
+    if (const MDNode *DeviceHasMD = F->getMetadata("sycl_declared_aspects")) {
+      // Entry points will have their declared aspects from their kernel call.
+      // To avoid double warnings, we skip them.
+      if (is_contained(EntryPoints, F))
+        continue;
+      for (const MDOperand &DeviceHasMDOp : DeviceHasMD->operands()) {
+        const auto *CAM = cast<ConstantAsMetadata>(DeviceHasMDOp);
+        const Constant *C = CAM->getValue();
+        DeviceHasAspectSet.insert(cast<ConstantInt>(C)->getSExtValue());
+      }
+      OriginatedFromAttribute = true;
+    } else if (F->hasFnAttribute("sycl-device-has")) {
+      Attribute DeviceHasAttr = F->getFnAttribute("sycl-device-has");
+      SmallVector<StringRef, 4> AspectValStrs;
+      DeviceHasAttr.getValueAsString().split(
+          AspectValStrs, ',', /*MaxSplit=*/-1, /*KeepEmpty=*/false);
+      for (StringRef AspectValStr : AspectValStrs) {
+        int AspectVal = -1;
+        assert(!AspectValStr.getAsInteger(10, AspectVal) &&
+               "Aspect value in sycl-device-has is not an integer.");
+        DeviceHasAspectSet.insert(AspectVal);
+      }
+      OriginatedFromAttribute = false;
+    } else {
+      continue;
     }
 
     for (int Aspect : Aspects) {
@@ -338,9 +359,19 @@ void validateUsedAspectsForFunctions(const FunctionToAspectsMapTy &Map,
             [=](auto AspectIt) { return Aspect == AspectIt.second; });
         assert(AspectNameIt != AspectValues.end() &&
                "Used aspect is not part of the existing aspects");
+        // We may encounter an entry point when using the device_has property.
+        // In this case we act like the usage came from the first callee to
+        // avoid repeat warnings on the same line.
+        Function *AdjustedOriginF =
+            is_contained(EntryPoints, F)
+                ? findFirstAspectUsageCallee(CG.find(F)->second, Map, Aspect)
+                : F;
+        assert(AdjustedOriginF &&
+               "Adjusted function pointer for aspect usage is null");
         SmallVector<Function *, 8> CallChain =
-            getAspectUsageChain(F, Map, CG, Aspect);
-        diagnoseAspectsMismatch(F, CallChain, AspectNameIt->first);
+            getAspectUsageChain(AdjustedOriginF, Map, CG, Aspect);
+        diagnoseAspectsMismatch(AdjustedOriginF, CallChain, AspectNameIt->first,
+                                OriginatedFromAttribute);
       }
     }
   }