Merge branch 'intel:sycl' into pablo/graph_enqueue_functions

Author: reble

clang/lib/Driver/ToolChains/HIPUtility.cpp

@@ -324,7 +324,10 @@ void HIP::constructHIPFatbinCommand(Compilation &C, const JobAction &JA,
       Args.MakeArgString(std::string("-output=").append(Output));
   BundlerArgs.push_back(BundlerOutputArg);
 
-  addOffloadCompressArgs(Args, BundlerArgs);
+  // For SYCL, the compression is occurring during the wrapping step, so we do
+  // not want to do additional compression here.
+  if (!JA.isDeviceOffloading(Action::OFK_SYCL))
+    addOffloadCompressArgs(Args, BundlerArgs);
 
   const char *Bundler = Args.MakeArgString(
       T.getToolChain().GetProgramPath("clang-offload-bundler"));

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1780,6 +1780,16 @@ void SYCLToolChain::AddImpliedTargetArgs(const llvm::Triple &Triple,
     ArgStringList TargArgs;
     Args.AddAllArgValues(TargArgs, options::OPT_Xs, options::OPT_Xs_separate);
     Args.AddAllArgValues(TargArgs, options::OPT_Xsycl_backend);
+    // For -Xsycl-target-backend=<triple> the triple value is used to push
+    // specific options to the matching device compilation using that triple.
+    // Scrutinize this to make sure we are only checking the values needed
+    // for the current device compilation.
+    for (auto *A : Args) {
+      if (!A->getOption().matches(options::OPT_Xsycl_backend_EQ))
+        continue;
+      if (getDriver().MakeSYCLDeviceTriple(A->getValue()) == Triple)
+        TargArgs.push_back(A->getValue(1));
+    }
     // Check for any -device settings.
     std::string DevArg;
     if (IsJIT || Device == "pvc" || hasPVCDevice(TargArgs, DevArg)) {

clang/lib/Sema/SemaSYCL.cpp

@@ -3167,7 +3167,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler {
 //        // code
 //      }
 //
-//      [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
+//      [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const
 //      {
 //        // code
 //      }

clang/lib/Sema/SemaSYCLDeclAttr.cpp

@@ -132,6 +132,13 @@ void SemaSYCL::checkDeprecatedSYCLAttributeSpelling(const ParsedAttr &A,
     return;
   }
 
+  // Additionally, diagnose deprecated [[intel::reqd_sub_group_size]] spelling
+  if (A.getKind() == ParsedAttr::AT_IntelReqdSubGroupSize && A.getScopeName() &&
+      A.getScopeName()->isStr("intel")) {
+    diagnoseDeprecatedAttribute(A, "sycl", "reqd_sub_group_size");
+    return;
+  }
+
   // Diagnose SYCL 2020 spellings in later SYCL modes.
   if (getLangOpts().getSYCLVersion() >= LangOptions::SYCL_2020) {
     // All attributes in the cl vendor namespace are deprecated in favor of a

clang/test/CodeGenSYCL/kernel-op-calls.cpp

@@ -11,7 +11,7 @@ class Functor1 {
 public:
  Functor1(){}
 
- [[intel::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
+ [[sycl::reqd_sub_group_size(4)]] void operator()(sycl::id<1> id) const {}
 
  [[sycl::work_group_size_hint(1, 2, 3)]] void operator()(sycl::id<2> id) const {}
 

clang/test/CodeGenSYCL/reqd-sub-group-size.cpp

@@ -7,25 +7,25 @@ queue q;
 
 class Functor16 {
 public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(16)]] void operator()() const {}
 };
 
 template <int SIZE>
 class Functor2 {
 public:
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
 };
 
 template <int N>
-[[intel::reqd_sub_group_size(N)]] void func() {}
+[[sycl::reqd_sub_group_size(N)]] void func() {}
 
 int main() {
   q.submit([&](handler &h) {
     Functor16 f16;
     h.single_task<class kernel_name1>(f16);
 
     h.single_task<class kernel_name3>(
-        []() [[intel::reqd_sub_group_size(4)]]{});
+        []() [[sycl::reqd_sub_group_size(4)]]{});
 
     Functor2<2> f2;
     h.single_task<class kernel_name4>(f2);

clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp

@@ -7,12 +7,12 @@ queue q;
 
 class Functor {
 public:
-  [[intel::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(4), cl::reqd_work_group_size(32, 16, 16)]] void operator()() const {}
 };
 
 class Functor1 {
 public:
-  [[intel::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(2), sycl::reqd_work_group_size(64, 32, 32)]] void operator()() const {}
 };
 
 template <int SIZE, int SIZE1, int SIZE2>

clang/test/Driver/sycl-ftarget-register-alloc-mode-old-model.cpp

@@ -20,6 +20,14 @@
 // RUN:    -fsycl-targets=spir64_gen -Xs "-device pvc" %s 2>&1 \
 // RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc
 
+// RUN: %clang -### -fsycl --no-offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" %s 2>&1 \
+// RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc
+
+// RUN: %clang -### -fsycl --no-offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" %s 2>&1 \
+// RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc
+
 // RUN: %clang -### -fsycl --no-offload-new-driver \
 // RUN:    -fsycl-targets=spir64_gen -Xs "-device 0x0BD5" %s 2>&1 \
 // RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc

clang/test/Driver/sycl-ftarget-register-alloc-mode.cpp

@@ -20,6 +20,14 @@
 // RUN:    -fsycl-targets=spir64_gen -Xs "-device pvc" %s 2>&1 \
 // RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc
 
+// RUN: %clang -### -fsycl --offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -Xsycl-target-backend "-device pvc" %s 2>&1 \
+// RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc
+
+// RUN: %clang -### -fsycl --offload-new-driver \
+// RUN:    -fsycl-targets=spir64_gen -Xsycl-target-backend=spir64_gen "-device pvc" %s 2>&1 \
+// RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc
+
 // RUN: %clang -### -fsycl --offload-new-driver \
 // RUN:    -fsycl-targets=spir64_gen -Xs "-device 0x0BD5" %s 2>&1 \
 // RUN:   | FileCheck %if system-windows %{ -check-prefix=DEFAULT_AOT %} %else %{ -check-prefix=AUTO_AOT %} %s -DDEVICE=pvc

clang/test/Driver/sycl-offload-new-driver.c

@@ -187,7 +187,7 @@
 // RUN: %clangxx -fsycl -### -fsycl-targets=spir64_gen --offload-new-driver \
 // RUN:   -Xsycl-target-backend=spir64_gen "-device pvc,bdw" %s 2>&1 \
 // RUN:   | FileCheck -check-prefix COMMA_FILE %s
-// COMMA_FILE: clang-offload-packager{{.*}} "--image=file={{.*}}pvc@bdw{{.*}},triple=spir64_gen-unknown-unknown,arch=pvc,bdw,kind=sycl"
+// COMMA_FILE: clang-offload-packager{{.*}} "--image=file={{.*}}pvc@bdw{{.*}},triple=spir64_gen-unknown-unknown,arch=pvc,bdw,kind=sycl,compile-opts=-device_options pvc -ze-intel-enable-auto-large-GRF-mode"
 
 /// Verify that --cuda-path is passed to clang-linker-wrapper for SYCL offload
 // RUN: %clangxx -fsycl -### -fsycl-targets=nvptx64-nvidia-cuda \