[SYCL] Adjust kernel parameters requirements

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10314,8 +10314,10 @@ def err_sycl_virtual_types : Error<
   "No class with a vtable can be used in a SYCL kernel or any code included in the kernel">;
 def note_sycl_used_here : Note<"used here">;
 def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">;
+def err_sycl_non_trivially_copyable_type : Error<
+  "kernel parameter has non-trivially copyable class/struct type %0">;
 def err_sycl_non_std_layout_type : Error<
-  "kernel parameter has non-standard layout class/struct type">;
+  "kernel parameter has non-standard layout class/struct type %0">;
 def err_conflicting_sycl_kernel_attributes : Error<
   "conflicting attributes applied to a SYCL kernel">;
 def err_conflicting_sycl_function_attributes : Error<

clang/include/clang/Basic/LangOptions.def

@@ -233,6 +233,7 @@ LANGOPT(GPUMaxThreadsPerBlock, 32, 256, "default max threads per block for kerne
 LANGOPT(SYCLIsDevice      , 1, 0, "Generate code for SYCL device")
 LANGOPT(SYCLIsHost        , 1, 0, "SYCL host compilation")
 LANGOPT(SYCLAllowFuncPtr  , 1, 0, "Allow function pointers in SYCL device code")
+LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters")
 LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")
 
 LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")

clang/include/clang/Driver/CC1Options.td

@@ -899,6 +899,8 @@ def fsycl_int_header : Separate<["-"], "fsycl-int-header">,
   HelpText<"Generate SYCL integration header into this file.">;
 def fsycl_int_header_EQ : Joined<["-"], "fsycl-int-header=">,
   Alias<fsycl_int_header>;
+def fsycl_std_layout_kernel_params: Flag<["-"], "fsycl-std-layout-kernel-params">,
+  HelpText<"Enable standard layout requirement for SYCL kernel parameters.">;
 def fsycl_allow_func_ptr : Flag<["-"], "fsycl-allow-func-ptr">,
   HelpText<"Allow function pointers in SYCL device.">;
 def fno_sycl_allow_func_ptr : Flag<["-"], "fno-sycl-allow-func-ptr">;

clang/lib/AST/RecordLayoutBuilder.cpp

@@ -2199,8 +2199,9 @@ static bool mustSkipTailPadding(TargetCXXABI ABI, const CXXRecordDecl *RD) {
   llvm_unreachable("bad tail-padding use kind");
 }
 
-static bool isMsLayout(const ASTContext &Context) {
-  return Context.getTargetInfo().getCXXABI().isMicrosoft();
+static bool isMsLayout(const ASTContext &Context, bool CheckAuxABI = false) {
+  return (CheckAuxABI) ? Context.getAuxTargetInfo()->getCXXABI().isMicrosoft()
+                       : Context.getTargetInfo().getCXXABI().isMicrosoft();
 }
 
 // This section contains an implementation of struct layout that is, up to the
@@ -3025,6 +3026,9 @@ ASTContext::getASTRecordLayout(const RecordDecl *D) const {
   // as soon as we begin to parse the definition.  That definition is
   // not a complete definition (which is what isDefinition() tests)
   // until we *finish* parsing the definition.
+  bool CheckAuxABI = false;
+  if (getLangOpts().SYCLIsDevice && (getAuxTargetInfo() != nullptr))
+    CheckAuxABI = true;
 
   if (D->hasExternalLexicalStorage() && !D->getDefinition())
     getExternalSource()->CompleteType(const_cast<RecordDecl*>(D));
@@ -3042,7 +3046,7 @@ ASTContext::getASTRecordLayout(const RecordDecl *D) const {
 
   const ASTRecordLayout *NewEntry = nullptr;
 
-  if (isMsLayout(*this)) {
+  if (isMsLayout(*this, CheckAuxABI)) {
     MicrosoftRecordLayoutBuilder Builder(*this);
     if (const auto *RD = dyn_cast<CXXRecordDecl>(D)) {
       Builder.cxxLayout(RD);

clang/lib/Driver/ToolChains/Clang.cpp

@@ -3989,6 +3989,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
 
   if (Arg *A = Args.getLastArg(options::OPT_sycl_std_EQ)) {
     A->render(Args, CmdArgs);
+    CmdArgs.push_back("-fsycl-std-layout-kernel-params");
   } else if (IsSYCL) {
     // Ensure the default version in SYCL mode is 1.2.1
     CmdArgs.push_back("-sycl-std=1.2.1");

clang/lib/Frontend/CompilerInvocation.cpp

@@ -3130,6 +3130,8 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
   Opts.SYCLIsHost   = Args.hasArg(options::OPT_fsycl_is_host);
   Opts.SYCLAllowFuncPtr = Args.hasFlag(options::OPT_fsycl_allow_func_ptr,
                                   options::OPT_fno_sycl_allow_func_ptr, false);
+  Opts.SYCLStdLayoutKernelParams =
+      Args.hasArg(options::OPT_fsycl_std_layout_kernel_params);
   Opts.SYCLUnnamedLambda = Args.hasArg(options::OPT_fsycl_unnamed_lambda);
 
   // Set CUDA mode for OpenMP target NVPTX if specified in options

clang/lib/Sema/SemaSYCL.cpp

@@ -969,7 +969,6 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) {
 // Returns true if all arguments are successfully built.
 static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
                         SmallVectorImpl<ParamDesc> &ParamDescs) {
-  const LambdaCapture *Cpt = KernelObj->captures_begin();
   auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) {
     // Create a parameter descriptor and append it to the result
     ParamDescs.push_back(makeParamDesc(Fld, ArgType));
@@ -1033,26 +1032,35 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
     QualType ArgTy = Fld->getType();
     if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) {
       createSpecialSYCLObjParamDesc(Fld, ArgTy);
-    } else if (!ArgTy->isStandardLayoutType()) {
-      // SYCL v1.2.1 s4.8.10 p5:
-      // C++ non-standard layout values must not be passed as arguments to a
-      // kernel that is compiled for a device.
-      const auto &DiagLocation =
-          Cpt ? Cpt->getLocation() : cast<DeclaratorDecl>(Fld)->getLocation();
-
-      Context.getDiagnostics().Report(DiagLocation,
-                                      diag::err_sycl_non_std_layout_type);
-
-      // Set the flag and continue processing so we can emit error for each
-      // invalid argument.
-      AllArgsAreValid = false;
     } else if (ArgTy->isStructureOrClassType()) {
-      assert(ArgTy->isStandardLayoutType());
+      if (Context.getLangOpts().SYCLStdLayoutKernelParams) {
+        if (!ArgTy->isStandardLayoutType()) {
+          Context.getDiagnostics().Report(Fld->getLocation(),
+                                          diag::err_sycl_non_std_layout_type)
+              << ArgTy;
+          AllArgsAreValid = false;
+          continue;
+        }
+      }
+      // TODO: Make stream class trivially copyable and remove the check on
+      // stream class.
+      if (!ArgTy.isTriviallyCopyableType(Context) &&
+          !Util::isSyclStreamType(ArgTy)) {
+        Context.getDiagnostics().Report(
+            Fld->getLocation(), diag::err_sycl_non_trivially_copyable_type)
+            << ArgTy;
+        AllArgsAreValid = false;
+        continue;
+      }
 
       CreateAndAddPrmDsc(Fld, ArgTy);
 
       // Create descriptors for each accessor field in the class or struct
       createParamDescForWrappedAccessors(Fld, ArgTy);
+    } else if (ArgTy->isReferenceType()) {
+      Context.getDiagnostics().Report(
+          Fld->getLocation(), diag::err_bad_kernel_param_type) << ArgTy;
+      AllArgsAreValid = false;
     } else if (ArgTy->isPointerType()) {
       // Pointer Arguments need to be in the global address space
       QualType PointeeTy = ArgTy->getPointeeType();
@@ -1068,10 +1076,6 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
     } else {
       llvm_unreachable("Unsupported kernel parameter type");
     }
-
-    // Update capture iterator as we process arguments
-    if (Cpt && Cpt != KernelObj->captures_end())
-      ++Cpt;
   }
 
   return AllArgsAreValid;

clang/test/CodeGenSYCL/non-standard-layout.cpp

@@ -0,0 +1,21 @@
+// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-unknown-windows-unknown -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck --check-prefix CHK-WIN %s
+// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-unknown-sycldevice -aux-triple x86_64-unknown-linux-unknown -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck --check-prefix CHK-LIN %s
+
+#include "sycl.hpp"
+// CHK-WIN: %struct{{.*}}F = type { i8, i8 }
+// CHK-LIN: %struct{{.*}}F = type { i8 }
+struct F1 {};
+struct F2 {};
+struct F : F1, F2 {
+  char x;
+};
+
+int main() {
+  cl::sycl::accessor<F, 1, cl::sycl::access::mode::read_write> accessorA;
+  cl::sycl::handler cgh;
+  cgh.single_task<class kernel_function>(
+      [=]() {
+        accessorA.use();
+      });
+  return 0;
+}

clang/test/SemaSYCL/non-std-layout-param.cpp

@@ -1,4 +1,5 @@
-// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-std-layout-kernel-params -verify -fsyntax-only %s
+// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only %s
 
 // This test checks if compiler reports compilation error on an attempt to pass
 // non-standard layout struct object as SYCL kernel parameter.
@@ -23,38 +24,11 @@ void test() {
   C C0;
   C0.Y=0;
   kernel_single_task<class MyKernel>([=] {
-    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
     (void)C0.Y;
   });
 }
 
-void test_capture_explicit_ref() {
-  int p = 0;
-  double q = 0;
-  float s = 0;
-  kernel_single_task<class kernel_capture_single_ref>([
-    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
-      &p,
-      q,
-    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
-      &s] {
-    (void) q;
-    (void) p;
-    (void) s;
-  });
-}
-
-void test_capture_implicit_refs() {
-  int p = 0;
-  double q = 0;
-  kernel_single_task<class kernel_capture_refs>([&] {
-    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
-    (void) p;
-    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
-    (void) q;
-  });
-}
-
 struct Kernel {
   void operator()() {
     (void) c1;
@@ -64,12 +38,12 @@ struct Kernel {
   }
 
   int p;
-  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
   C c1;
 
   int q;
 
-  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
+  // expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
   C c2;
 };
 

clang/test/SemaSYCL/non-trivially-copyable-kernel-param.cpp

@@ -0,0 +1,28 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+
+// This test checks if compiler reports compilation error on an attempt to pass
+// a struct with non-trivially copyable type as SYCL kernel parameter.
+
+struct A { int i; };
+
+struct B {
+  int i;
+  B (int _i) : i(_i) {}
+  B (const B& x) : i(x.i) {}
+};
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+void test() {
+  A IamGood;
+  IamGood.i = 0;
+  B IamBad(1);
+  kernel_single_task<class kernel_capture_refs>([=] {
+    int a = IamGood.i;
+    // expected-error@+1 {{kernel parameter has non-trivially copyable class/struct type}}
+    int b = IamBad.i;
+  });
+}

clang/test/SemaSYCL/reference-kernel-param.cpp

@@ -0,0 +1,36 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
+
+// This test checks if compiler reports compilation error on an attempt to pass
+// a reference as SYCL kernel parameter.
+
+template <typename Name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+void test_capture_explicit_ref() {
+  int p = 0;
+  double q = 0;
+  float s = 0;
+  kernel_single_task<class kernel_capture_single_ref>([
+    // expected-error@+1 {{'int &' cannot be used as the type of a kernel parameter}}
+      &p,
+      q,
+    // expected-error@+1 {{'float &' cannot be used as the type of a kernel parameter}}
+      &s] {
+    (void) q;
+    (void) p;
+    (void) s;
+  });
+}
+
+void test_capture_implicit_refs() {
+  int p = 0;
+  double q = 0;
+  kernel_single_task<class kernel_capture_refs>([&] {
+    // expected-error@+1 {{'int &' cannot be used as the type of a kernel parameter}}
+    (void) p;
+    // expected-error@+1 {{'double &' cannot be used as the type of a kernel parameter}}
+    (void) q;
+  });
+}

sycl/test/basic_tests/handler/handler_mem_op.cpp

@@ -19,7 +19,7 @@
 using namespace cl::sycl;
 
 template <typename T> struct point {
-  point(const point &rhs) : x(rhs.x), y(rhs.y) {}
+  point(const point &rhs) = default;
   point(T x, T y) : x(x), y(y) {}
   point(T v) : x(v), y(v) {}
   point() : x(0), y(0) {}

sycl/test/struct_param/non-standard-layout.cpp

@@ -0,0 +1,45 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+struct F1 {};
+struct F2 {};
+struct F : F1, F2 {
+  cl::sycl::cl_char x;
+};
+
+bool test0() {
+  F S;
+  S.x = 0;
+  F S0;
+  S0.x = 1;
+  {
+    buffer<F, 1> Buf(&S0, range<1>(1));
+    queue myQueue;
+    myQueue.submit([&](handler &cgh) {
+      auto B = Buf.get_access<access::mode::write>(cgh);
+      cgh.single_task<class NonStandard>([=] { B[0] = S; });
+    });
+  }
+  bool Passed = (S.x == S0.x);
+
+  if (!Passed) {
+    std::cout << "test0 failed" << std::endl;
+  }
+  return Passed;
+}
+
+int main() {
+
+  bool Pass = test0();
+
+  std::cout << "Test " << (Pass ? "passed" : "FAILED") << std::endl;
+  return Pass ? 0 : 1;
+
+}