Merge branch 'sycl' into fabio/binary_update_fix

Author: aarongreig

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/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-offload-wrapper-compression.cpp

@@ -10,5 +10,11 @@
 // RUN: %clangxx -### -fsycl %s 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
 // RUN: %clangxx -### -fsycl --offload-compression-level=3 %s 2>&1 | FileCheck %s --check-prefix=CHECK-NO-COMPRESS
 
+// For SYCL offloading to HIP, make sure we don't pass '--compress' to offload-bundler.
+// RUN: %clangxx -### -fsycl --offload-compress -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx1031 -fsycl-targets=amdgcn-amd-amdhsa -fno-sycl-libspirv -nogpulib %s &> %t.driver
+// RUN: FileCheck %s --check-prefix=CHECK-NO-COMPRESS-BUNDLER --input-file=%t.driver
+
 // CHECK-NO-COMPRESS-NOT: {{.*}}clang-offload-wrapper{{.*}}"-offload-compress"{{.*}}
 // CHECK-NO-COMPRESS-NOT: {{.*}}clang-offload-wrapper{{.*}}"-offload-compression-level=3"{{.*}}
+
+// CHECK-NO-COMPRESS-BUNDLER-NOT: {{.*}}clang-offload-bundler{{.*}}"-compress"{{.*}}

clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp

@@ -6,7 +6,7 @@ template <typename T> class Fobj {
 public:
   Fobj() {}
   void operator()() const {
-    auto L0 = []() [[intel::reqd_sub_group_size(4)]]{};
+    auto L0 = []() [[sycl::reqd_sub_group_size(4)]]{};
     L0();
   }
 };

clang/test/SemaSYCL/reqd-sub-group-size-ast.cpp

@@ -9,14 +9,14 @@ queue q;
 
 class Functor16 {
 public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(16)]] void operator()() const {}
 };
 
 // Test that checks template parameter support on member function of class template.
 template <int SIZE>
 class KernelFunctor {
 public:
-  [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {}
+  [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {}
 };
 
 // Test that checks template parameter support on function.
@@ -35,7 +35,7 @@ class KernelFunctor {
 // CHECK-NEXT: NonTypeTemplateParmDecl
 // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 12
 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) {
@@ -52,14 +52,14 @@ int main() {
     // CHECK-NEXT: ConstantExpr {{.*}} 'int'
     // CHECK-NEXT: value: Int 2
     // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}}
-    h.single_task<class kernel_name3>([]() [[intel::reqd_sub_group_size(2)]] {});
+    h.single_task<class kernel_name3>([]() [[sycl::reqd_sub_group_size(2)]] {});
 
     // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5
     // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
     // CHECK-NEXT: ConstantExpr {{.*}} 'int'
     // CHECK-NEXT: value: Int 6
     // CHECK-NEXT: IntegerLiteral{{.*}}6{{$}}
-    h.single_task<class kernel_name5>([]() [[intel::reqd_sub_group_size(6)]] {});
+    h.single_task<class kernel_name5>([]() [[sycl::reqd_sub_group_size(6)]] {});
 
     // CHECK: FunctionDecl {{.*}}kernel_name_6
     // CHECK: IntelReqdSubGroupSizeAttr {{.*}} reqd_sub_group_size
@@ -79,8 +79,8 @@ int main() {
         // CHECK-NEXT: value: Int 8
         // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}
         // CHECK-NOT: IntelReqdSubGroupSizeAttr
-        []() [[intel::reqd_sub_group_size(8),
-               intel::reqd_sub_group_size(8)]] {});
+        []() [[sycl::reqd_sub_group_size(8),
+               sycl::reqd_sub_group_size(8)]] {});
   });
   func<12>();
   return 0;

clang/test/SemaSYCL/reqd-sub-group-size-host.cpp

@@ -1,9 +1,9 @@
 // RUN: %clang_cc1 -fsycl-is-host -fsyntax-only -verify %s
 // expected-no-diagnostics
 
-[[intel::reqd_sub_group_size(8)]] void fun() {}
+[[sycl::reqd_sub_group_size(8)]] void fun() {}
 
 class Functor {
 public:
-  [[intel::reqd_sub_group_size(16)]] void operator()() {}
+  [[sycl::reqd_sub_group_size(16)]] void operator()() {}
 };