Merge from 'sycl' to 'sycl-web'

Author: iclsrc

.github/workflows/sycl_post_commit.yml

@@ -9,6 +9,10 @@ on:
     - sycl
     paths:
     - .github/workflows/sycl_post_commit.yml
+    - .github/workflows/sycl_gen_test_matrix.yml
+    - .github/workflows/sycl_linux_build_and_test.yml
+    - .github/workflows/sycl_windows_build_and_test.yml
+    - .github/workflows/sycl_macos_build_and_test.yml
 
 jobs:
   # This job generates matrix of tests for LLVM Test Suite

.github/workflows/sycl_precommit.yml

@@ -9,6 +9,8 @@ on:
     - '.github/ISSUE_TEMPLATE/**'
     - '.github/CODEOWNERS'
     - '.github/workflows/sycl_update_gpu_driver.yml'
+    - '.github/workflows/sycl_windows_build_and_test.yml'
+    - '.github/workflows/sycl_macos_build_and_test.yml'
     - 'devops/containers/**'
     - 'devops/scripts/install_drivers.sh'
     - 'devops/scripts/install_build_tools.sh'

.github/workflows/sycl_windows_build_and_test.yml

@@ -14,7 +14,7 @@ jobs:
     runs-on: [Windows, build]
     # TODO use cached checkout
     steps:
-    - uses: ilammy/msvc-dev-cmd@9f8ae839b01883414208f29e3e24524387f48e1f
+    - uses: ilammy/msvc-dev-cmd@7315a94840631165970262a99c72cfb48a65d25d
       with:
         arch: amd64
     - name: Set env

clang/lib/Sema/SemaSYCL.cpp

@@ -988,6 +988,28 @@ static QualType GetSYCLKernelObjectType(const FunctionDecl *KernelCaller) {
   return KernelParamTy.getUnqualifiedType();
 }
 
+static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
+  for (auto *MD : Rec->methods()) {
+    if (MD->getOverloadedOperator() == OO_Call)
+      return MD;
+  }
+  return nullptr;
+}
+
+// Fetch the associated call operator of the kernel object
+// (of either the lambda or the function object).
+static CXXMethodDecl *
+GetCallOperatorOfKernelObject(const CXXRecordDecl *KernelObjType) {
+  CXXMethodDecl *CallOperator = nullptr;
+  if (!KernelObjType)
+    return CallOperator;
+  if (KernelObjType->isLambda())
+    CallOperator = KernelObjType->getLambdaCallOperator();
+  else
+    CallOperator = getOperatorParens(KernelObjType);
+  return CallOperator;
+}
+
 /// Creates a kernel parameter descriptor
 /// \param Src  field declaration to construct name from
 /// \param Ty   the desired parameter type
@@ -2775,14 +2797,6 @@ class SyclOptReportCreator : public SyclKernelFieldHandler {
   }
 };
 
-static CXXMethodDecl *getOperatorParens(const CXXRecordDecl *Rec) {
-  for (auto *MD : Rec->methods()) {
-    if (MD->getOverloadedOperator() == OO_Call)
-      return MD;
-  }
-  return nullptr;
-}
-
 static bool isESIMDKernelType(const CXXRecordDecl *KernelObjType) {
   const CXXMethodDecl *OpParens = getOperatorParens(KernelObjType);
   return (OpParens != nullptr) && OpParens->hasAttr<SYCLSimdAttr>();
@@ -2871,13 +2885,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 
     // Fetch the kernel object and the associated call operator
     // (of either the lambda or the function object).
-    CXXRecordDecl *KernelObj =
+    const CXXRecordDecl *KernelObj =
         GetSYCLKernelObjectType(KernelCallerFunc)->getAsCXXRecordDecl();
-    CXXMethodDecl *WGLambdaFn = nullptr;
-    if (KernelObj->isLambda())
-      WGLambdaFn = KernelObj->getLambdaCallOperator();
-    else
-      WGLambdaFn = getOperatorParens(KernelObj);
+    CXXMethodDecl *WGLambdaFn = GetCallOperatorOfKernelObject(KernelObj);
+
     assert(WGLambdaFn && "non callable object is passed as kernel obj");
     // Mark the function that it "works" in a work group scope:
     // NOTE: In case of parallel_for_work_item the marker call itself is
@@ -3534,7 +3545,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
 static bool IsSYCLUnnamedKernel(Sema &SemaRef, const FunctionDecl *FD) {
   if (!SemaRef.getLangOpts().SYCLUnnamedLambda)
     return false;
-  QualType FunctorTy = GetSYCLKernelObjectType(FD);
+  const QualType FunctorTy = GetSYCLKernelObjectType(FD);
   QualType TmplArgTy = calculateKernelNameType(SemaRef.Context, FD);
   return SemaRef.Context.hasSameType(FunctorTy, TmplArgTy);
 }
@@ -3960,7 +3971,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc,
   const CXXRecordDecl *KernelObj =
       GetSYCLKernelObjectType(KernelFunc)->getAsCXXRecordDecl();
 
-  if (!KernelObj) {
+  if (!GetCallOperatorOfKernelObject(KernelObj)) {
     Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object);
     KernelFunc->setInvalidDecl();
     return;

clang/test/SemaSYCL/undefined-functor.cpp

@@ -0,0 +1,40 @@
+// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -verify -fsyntax-only %s
+// This test checks that an error is thrown when a functor without a call operator defined is used as a kernel.
+
+#include "sycl.hpp"
+
+using namespace sycl;
+queue q;
+
+struct FunctorWithoutCallOperator; // expected-note {{forward declaration of 'FunctorWithoutCallOperator'}}
+
+struct StructDefined {
+  int x;
+};
+
+class FunctorWithCallOpDefined {
+  int x;
+  public:
+  void operator()() const {}
+};
+
+int main() {
+  
+  q.submit([&](sycl::handler &cgh) {
+    // expected-error@#KernelSingleTask {{kernel parameter must be a lambda or function object}}
+    // expected-error@+2 {{invalid use of incomplete type 'FunctorWithoutCallOperator'}}
+    // expected-note@+1 {{in instantiation of function template specialization}}
+    cgh.single_task(FunctorWithoutCallOperator{});
+  });
+
+  q.submit([&](sycl::handler &cgh) {
+    // expected-error@#KernelSingleTask {{kernel parameter must be a lambda or function object}}
+    // expected-note@+1 {{in instantiation of function template specialization}}
+    cgh.single_task(StructDefined{});
+  });
+  
+  q.submit([&](sycl::handler &cgh) {
+    cgh.single_task(FunctorWithCallOpDefined{});
+  });
+
+}

sycl/include/sycl/accessor.hpp

@@ -1949,7 +1949,12 @@ class __SYCL_SPECIAL_CLASS __SYCL_TYPE(accessor) accessor :
 #endif
   }
 
-  void swap(accessor &other) { std::swap(impl, other.impl); }
+  void swap(accessor &other) {
+    std::swap(impl, other.impl);
+#ifndef __SYCL_DEVICE_ONLY__
+    std::swap(MAccData, other.MAccData);
+#endif
+  }
 
   constexpr bool is_placeholder() const { return IsPlaceH; }