[SYCL] Add SYCLPropagateAspectUsage Pass

clang/include/clang/Basic/DiagnosticFrontendKinds.td

@@ -313,4 +313,11 @@ def warn_profile_data_unprofiled : Warning<
   InGroup<ProfileInstrUnprofiled>;
 } // end of instrumentation issue category
 
+def warn_sycl_wrong_aspect_usage : Warning<
+  "function '%0' uses aspect '%1' not listed in 'sycl::device_has()'">,
+  InGroup<DiagGroup<"sycl-aspect-usage">>;
+def note_sycl_wrong_aspect_usage : Note<
+  "%select{the actual use is in|which is called by}0 "
+  "%1%select{|, compile with '-g' to get source location}2">;
+
 }

clang/lib/CodeGen/BackendUtil.cpp

@@ -45,6 +45,7 @@
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
+#include "llvm/SYCLLowerIR/PropagateAspectUsage.h"
 #include "llvm/Support/BuryPointer.h"
 #include "llvm/Support/CommandLine.h"
 #include "llvm/Support/MemoryBuffer.h"
@@ -1024,6 +1025,9 @@ void EmitAssemblyHelper::EmitAssemblyWithLegacyPassManager(
   PerFunctionPasses.add(
       createTargetTransformInfoWrapperPass(getTargetIRAnalysis()));
 
+  if (LangOpts.SYCLIsDevice)
+    PerModulePasses.add(createPropagateAspectUsagePass());
+
   CreatePasses(PerModulePasses, PerFunctionPasses);
 
   // Add a verifier pass if requested. We don't have to do this if the action
@@ -1382,6 +1386,9 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
 
   ModulePassManager MPM;
 
+  if (CodeGenOpts.DisableLLVMPasses && LangOpts.SYCLIsDevice)
+    MPM.addPass(PropagateAspectUsagePass());
+
   if (!CodeGenOpts.DisableLLVMPasses) {
     // Map our optimization levels into one of the distinct levels used to
     // configure the pipeline.
@@ -1391,6 +1398,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
       PB.registerPipelineStartEPCallback(
           [](ModulePassManager &MPM, OptimizationLevel Level) {
             MPM.addPass(ESIMDVerifierPass());
+            MPM.addPass(PropagateAspectUsagePass());
           });
 
     bool IsThinLTO = CodeGenOpts.PrepareForThinLTO;

clang/lib/CodeGen/CodeGenAction.cpp

@@ -444,6 +444,7 @@ namespace clang {
     void OptimizationFailureHandler(
         const llvm::DiagnosticInfoOptimizationFailure &D);
     void DontCallDiagHandler(const DiagnosticInfoDontCall &D);
+    void SYCLWarningDiagHandler(const DiagnosticInfoSYCLUnspecAspect &D);
   };
 
   void BackendConsumer::anchor() {}
@@ -825,6 +826,27 @@ void BackendConsumer::DontCallDiagHandler(const DiagnosticInfoDontCall &D) {
       << llvm::demangle(D.getFunctionName().str()) << D.getNote();
 }
 
+void BackendConsumer::SYCLWarningDiagHandler(
+    const DiagnosticInfoSYCLUnspecAspect &D) {
+  // FIXME: instead of querying source location from debug info, we should
+  // re-use existing mechanism with srcloc metadata, see D106030, D110364.
+  SourceLocation Loc;
+  Diags.Report(Loc, diag::warn_sycl_wrong_aspect_usage)
+      << D.getFunctionName() << D.getAspect();
+
+  if (!D.getCallChain().empty()) {
+    Diags.Report(Loc, diag::note_sycl_wrong_aspect_usage)
+        << /* the actual use */ 0 << D.getCallChain().back()
+        << !D.isFullDebugMode();
+
+    for (auto I = ++D.getCallChain().rbegin(), E = D.getCallChain().rend();
+         I != E; ++I) {
+      Diags.Report(Loc, diag::note_sycl_wrong_aspect_usage)
+          << /* called by */ 1 << *I << !D.isFullDebugMode();
+    }
+  }
+}
+
 /// This function is invoked when the backend needs
 /// to report something to the user.
 void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) {
@@ -899,6 +921,9 @@ void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) {
   case llvm::DK_DontCall:
     DontCallDiagHandler(cast<DiagnosticInfoDontCall>(DI));
     return;
+  case llvm::DK_SYCLWarning:
+    SYCLWarningDiagHandler(cast<DiagnosticInfoSYCLUnspecAspect>(DI));
+    return;
   default:
     // Plugin IDs are not bound to any value as they are set dynamically.
     ComputeDiagRemarkID(Severity, backend_plugin, DiagID);

clang/test/CodeGenSYCL/aspect_usage.cpp

@@ -0,0 +1,20 @@
+// Checks that Propagate aspect usage Pass is run for SYCL device target.
+//
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown %s -flegacy-pass-manager -mllvm -debug-pass=Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-SYCL1
+// CHECK-SYCL1: Propagate aspect usage
+//
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown %s -flegacy-pass-manager -fno-sycl-early-optimizations -mllvm -debug-pass=Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-SYCL2
+// CHECK-SYCL2: Propagate aspect usage
+//
+// RUN: %clang_cc1 %s -flegacy-pass-manager -mllvm -debug-pass=Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-NOT-SYCL1
+// CHECK-NOT-SYCL1-NOT: Propagate aspect usage
+//
+//
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown %s -fno-legacy-pass-manager -mdebug-pass Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-NEWPM-SYCL1
+// CHECK-NEWPM-SYCL1: PropagateAspectUsagePass
+//
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown %s -fno-legacy-pass-manager -fno-sycl-early-optimizations -mdebug-pass Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-NEWPM-SYCL2
+// CHECK-NEWPM-SYCL2: PropagateAspectUsagePass
+//
+// RUN: %clang_cc1 %s -fno-legacy-pass-manager -mdebug-pass Structure -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix=CHECK-NEWPM-NOT-SYCL
+// CHECK-NEWPM-NOT-SYCL-NOT: PropagateAspectUsagePass

clang/test/CodeGenSYCL/uses_aspects.cpp

@@ -28,7 +28,7 @@ template <cl::sycl::aspect Aspect>
 void func5() {}
 
 [[__sycl_detail__::__uses_aspects__(cl::sycl::aspect::cpu)]] void func6();
-// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !intel_used_aspects ![[ASPECTS1]] {
+// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !intel_used_aspects ![[ASPECTS4:[0-9]+]] {
 void func6() {
   Type1WithAspect TestObj1;
   Type2WithAspect TestObj2;
@@ -57,10 +57,11 @@ void foo() {
     h.single_task<class kernel_name_1>(f1);
   });
 }
-// CHECK: !intel_types_that_use_aspects = !{![[TYPE1:[0-9]+]], ![[TYPE2:[0-9]+]]}
-// CHECK: [[TYPE1]] = !{!"class.Type1WithAspect", i32 1}
-// CHECK: [[TYPE2]] = !{!"class.Type2WithAspect", i32 5, i32 1}
-// CHECK: [[EMPTYASPECTS]] = !{}
-// CHECK: [[ASPECTS1]] = !{i32 1}
-// CHECK: [[ASPECTS2]] = !{i32 5, i32 2}
-// CHECK: [[ASPECTS3]] = !{i32 0}
+// CHECK-DAG: !intel_types_that_use_aspects = !{![[TYPE1:[0-9]+]], ![[TYPE2:[0-9]+]]}
+// CHECK-DAG: [[TYPE1]] = !{!"class.Type1WithAspect", i32 1}
+// CHECK-DAG: [[TYPE2]] = !{!"class.Type2WithAspect", i32 5, i32 1}
+// CHECK-DAG: [[EMPTYASPECTS]] = !{}
+// CHECK-DAG: [[ASPECTS1]] = !{i32 1}
+// CHECK-DAG: [[ASPECTS2]] = !{i32 5, i32 2}
+// CHECK-DAG: [[ASPECTS3]] = !{i32 0}
+// CHECK-DAG: [[ASPECTS4]] = !{i32 1, i32 5}

clang/test/Frontend/Inputs/sycl.hpp

@@ -0,0 +1,68 @@
+#pragma once
+
+#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+
+inline namespace cl {
+namespace sycl {
+
+// Dummy aspect enum with limited enumerators
+enum class aspect {
+  host = 0,
+  cpu = 1,
+  gpu = 2,
+  accelerator = 3,
+  custom = 4,
+  fp16 = 5,
+  fp64 = 6,
+};
+
+class kernel {};
+class context {};
+class device {};
+class event {};
+
+class queue {
+public:
+  template <typename T>
+  event submit(T cgf) { return event{}; }
+
+  void wait() {}
+  void wait_and_throw() {}
+  void throw_asynchronous() {}
+};
+
+class auto_name {};
+template <typename Name, typename Type>
+struct get_kernel_name_t {
+  using name = Name;
+};
+template <typename Type>
+struct get_kernel_name_t<auto_name, Type> {
+  using name = Type;
+};
+
+class kernel_handler {
+  void __init_specialization_constants_buffer(char *specialization_constants_buffer) {}
+};
+
+#define ATTR_SYCL_KERNEL __attribute__((sycl_kernel))
+template <typename KernelName, typename KernelType>
+ATTR_SYCL_KERNEL void kernel_single_task(const KernelType &kernelFunc) { // #KernelSingleTask
+  kernelFunc();
+}
+
+class handler {
+public:
+  template <typename KernelName = auto_name, typename KernelType>
+  void single_task(const KernelType &kernelFunc) {
+    using NameT = typename get_kernel_name_t<KernelName, KernelType>::name;
+#ifdef __SYCL_DEVICE_ONLY__
+    kernel_single_task<NameT>(kernelFunc);
+#else
+    kernelFunc();
+#endif
+  }
+};
+
+} // namespace sycl
+} // namespace cl

clang/test/Frontend/sycl-aspects-usage.cpp

@@ -0,0 +1,107 @@
+// Test checks Propagate Aspect Usage pass.
+
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-WARNINGS
+
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown %s -debug-info-kind=constructor -dwarf-version=5 -emit-llvm -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-WARNINGS-DBG -DPATH=%s
+
+// CHECK-WARNINGS: warning: function 'checkStructUsesAspect(int)' uses aspect 'fp16' not listed in 'sycl::device_has()'
+// CHECK-WARNINGS-NEXT: note: the actual use is in funcWithStruct(int), compile with '-g' to get source location
+// CHECK-WARNINGS-NEXT: note: which is called by checkStructUsesAspect(int), compile with '-g' to get source location
+//
+// CHECK-WARNINGS-DBG: warning: function 'checkStructUsesAspect(int)' uses aspect 'fp16' not listed in 'sycl::device_has()'
+// CHECK-WARNINGS-DBG-NEXT: note: the actual use is in funcWithStruct(int) at [[PATH]]:27:5
+// CHECK-WARNINGS-DBG-NEXT: note: which is called by checkStructUsesAspect(int) at [[PATH]]:32:10
+
+#include "Inputs/sycl.hpp"
+
+using namespace cl::sycl;
+
+class KernelName;
+
+struct [[__sycl_detail__::__uses_aspects__(aspect::fp16)]] Struct {
+  int a = 0;
+};
+
+int funcWithStruct(int) {
+  Struct s;
+  s.a = 1;
+  return s.a;
+}
+
+[[sycl::device_has(aspect::fp64)]] int checkStructUsesAspect(int) {
+  return funcWithStruct(1);
+}
+
+// Check that empty device_has() emits a warning.
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-EMPTY
+// CHECK-EMPTY: warning: function 'checkEmptyDeviceHas()' uses aspect 'fp16' not listed in 'sycl::device_has()'
+[[sycl::device_has()]] int checkEmptyDeviceHas() {
+  return funcWithStruct(1);
+}
+
+[[sycl::device_has(aspect::fp16)]] int func2() {
+  return funcWithStruct(1);
+}
+
+// Check that empty device_has() emits a warning despite the fact
+// that invoked function's device_has() attribute is conformant
+// with actual usage.
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-EMPTY2
+// CHECK-EMPTY2: warning: function 'checkEmptyDeviceHas()' uses aspect 'fp16' not listed in 'sycl::device_has()'
+[[sycl::device_has()]] int checkEmptyDeviceHas2() {
+  return func2();
+}
+
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-DOUBLE
+// CHECK-DOUBLE: warning: function 'checkDouble()' uses aspect 'fp64' not listed in 'sycl::device_has()'
+// CHECK-DOUBLE-NEXT: note: the actual use is in checkDouble(), compile with '-g' to get source location
+
+[[sycl::device_has(aspect::fp16)]] int checkDouble() {
+  double d = 123;
+  // Strange calculations just to prevent AST optimizations
+  for (int i = 0; i < 10; ++i)
+    d += d * d;
+
+  return d;
+}
+
+struct [[__sycl_detail__::__uses_aspects__(aspect::cpu)]] StructWithCpu {
+  int a = 0;
+};
+
+int funcWithSeveralAspects() {
+  Struct s1;
+  StructWithCpu s2;
+  return static_cast<double>(s1.a);
+}
+
+// Check that a warning diagnostic works in a case
+// when there are several aspects present which conflict
+// with declared in device_has().
+
+// Check for fp64 aspect
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-SEVERAL-FP64
+// CHECK-SEVERAL-FP64: warning: function 'checkSeveralAspects()' uses aspect 'fp64' not listed in 'sycl::device_has()'
+// CHECK-SEVERAL-FP64-NEXT: note: the actual use is in funcWithSeveralAspects(), compile with '-g' to get source location
+// CHECK-SEVERAL-FP64-NEXT: note: which is called by checkSeveralAspects(), compile with '-g' to get source location
+//
+// Check for cpu aspect
+// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm %s -o /dev/null 2>&1 | FileCheck %s --check-prefix CHECK-SEVERAL-CPU
+// CHECK-SEVERAL-CPU: warning: function 'checkSeveralAspects()' uses aspect 'cpu' not listed in 'sycl::device_has()'
+// CHECK-SEVERAL-CPU-NEXT: note: the actual use is in funcWithSeveralAspects(), compile with '-g' to get source location
+// CHECK-SEVERAL-CPU-NEXT: note: which is called by checkSeveralAspects(), compile with '-g' to get source location
+[[sycl::device_has(aspect::fp16)]] int checkSeveralAspects() {
+  return funcWithSeveralAspects();
+}
+
+int main() {
+  queue Q;
+  Q.submit([&](handler &h) {
+    h.single_task<KernelName>([=]() {
+      checkStructUsesAspect(1);
+      checkEmptyDeviceHas();
+      checkDouble();
+      checkSeveralAspects();
+    });
+  });
+}

llvm/include/llvm/IR/DiagnosticInfo.h

@@ -17,6 +17,7 @@
 #include "llvm-c/Types.h"
 #include "llvm/ADT/ArrayRef.h"
 #include "llvm/ADT/Optional.h"
+#include "llvm/ADT/SmallString.h"
 #include "llvm/ADT/SmallVector.h"
 #include "llvm/ADT/StringRef.h"
 #include "llvm/ADT/Twine.h"
@@ -85,6 +86,7 @@ enum DiagnosticKind {
   DK_Unsupported,
   DK_SrcMgr,
   DK_DontCall,
+  DK_SYCLWarning,
   DK_FirstPluginKind // Must be last value to work with
                      // getNextAvailablePluginDiagnosticKind
 };
@@ -1098,6 +1100,36 @@ class DiagnosticInfoDontCall : public DiagnosticInfo {
   }
 };
 
+/// Class for emiting warning messages from SYCL PropagateAspectUsage pass
+/// about unspecified aspect.
+class DiagnosticInfoSYCLUnspecAspect : public DiagnosticInfo {
+  SmallString<16> FunctionName;
+  SmallString<5> AspectStr;
+  SmallVector<std::string, 8> CallChain;
+  bool IsFullDebugMode = false;
+
+public:
+  DiagnosticInfoSYCLUnspecAspect(StringRef FunctionName, StringRef AspectStr,
+                                 const SmallVector<std::string, 8> &CallChain,
+                                 bool IsFullDebugMode)
+      : DiagnosticInfo(DK_SYCLWarning, DiagnosticSeverity::DS_Warning),
+        FunctionName(FunctionName), AspectStr(AspectStr),
+        CallChain(std::move(CallChain)), IsFullDebugMode(IsFullDebugMode) {}
+
+  StringRef getFunctionName() const { return FunctionName; }
+
+  StringRef getAspect() const { return AspectStr; }
+
+  const SmallVector<std::string, 8> &getCallChain() const { return CallChain; }
+
+  int isFullDebugMode() const { return IsFullDebugMode; }
+
+  void print(DiagnosticPrinter &DP) const override;
+  static bool classof(const DiagnosticInfo *DI) {
+    return DI->getKind() == DK_SYCLWarning;
+  }
+};
+
 } // end namespace llvm
 
 #endif // LLVM_IR_DIAGNOSTICINFO_H

llvm/include/llvm/InitializePasses.h

@@ -448,6 +448,7 @@ void initializeESIMDLowerLoadStorePass(PassRegistry &);
 void initializeESIMDLowerVecArgLegacyPassPass(PassRegistry &);
 void initializeESIMDVerifierPass(PassRegistry &);
 void initializeSYCLLowerWGLocalMemoryLegacyPass(PassRegistry &);
+void initializeSYCLPropagateAspectUsageLegacyPassPass(PassRegistry &);
 void initializeTailCallElimPass(PassRegistry&);
 void initializeTailDuplicatePass(PassRegistry&);
 void initializeTargetLibraryInfoWrapperPassPass(PassRegistry&);

llvm/include/llvm/LinkAllPasses.h

@@ -39,6 +39,7 @@
 #include "llvm/IR/IRPrintingPasses.h"
 #include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
 #include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
+#include "llvm/SYCLLowerIR/PropagateAspectUsage.h"
 #include "llvm/Support/Valgrind.h"
 #include "llvm/Transforms/AggressiveInstCombine/AggressiveInstCombine.h"
 #include "llvm/Transforms/IPO.h"
@@ -209,6 +210,7 @@ namespace {
       (void) llvm::createExpandMemCmpPass();
       (void) llvm::createExpandVectorPredicationPass();
       (void)llvm::createESIMDVerifierPass();
+      (void)llvm::createPropagateAspectUsagePass();
       std::string buf;
       llvm::raw_string_ostream os(buf);
       (void) llvm::createPrintModulePass(os);