Apply code review comments

Signed-off-by: Mariya Podchishchaeva <[email protected]>

Author: Fznamznon

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10316,6 +10316,8 @@ 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 %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(SYCLNewKernelParamReq, 1, 0, "Enable trivially copyable requirement instead of 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_new_kernel_param_requirements: Flag<["-"], "fsycl-new-kernel-param-requirements">,
+  HelpText<"Enable trivially copyable requirement instead of 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/Driver/ToolChains/Clang.cpp

@@ -3992,6 +3992,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
   } else if (IsSYCL) {
     // Ensure the default version in SYCL mode is 1.2.1
     CmdArgs.push_back("-sycl-std=1.2.1");
+    CmdArgs.push_back("-fsycl-new-kernel-param-requirements");
   }
 
   if (IsOpenMPDevice) {

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.SYCLNewKernelParamReq =
+      Args.hasArg(options::OPT_fsycl_new_kernel_param_requirements);
   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

@@ -1033,15 +1033,25 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
     if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) {
       createSpecialSYCLObjParamDesc(Fld, ArgTy);
     } else if (ArgTy->isStructureOrClassType()) {
-      // 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;
+      if (Context.getLangOpts().SYCLNewKernelParamReq) {
+        // 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;
+        }
+      } else {
+        if (!ArgTy->isStandardLayoutType()) {
+          Context.getDiagnostics().Report(Fld->getLocation(),
+                                          diag::err_sycl_non_std_layout_type)
+              << ArgTy;
+          AllArgsAreValid = false;
+          continue;
+        }
       }
 
       CreateAndAddPrmDsc(Fld, ArgTy);

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

@@ -1,5 +1,5 @@
-// RUN: %clang_cc1 -I %S/Inputs -triple spir64-unknown-linux-sycldevice -aux-triple x86_64-unknown-windows-unknown -std=c++11 -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-linux-sycldevice -aux-triple x86_64-unknown-linux-unknown -std=c++11 -fsycl-is-device -disable-llvm-passes -S -emit-llvm %s -o - | FileCheck --check-prefix CHK-LIN %s
+// 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 }
@@ -10,14 +10,10 @@ struct F : F1, F2 {
   char x;
 };
 
-template <typename name, typename Func>
-__attribute__((sycl_kernel)) void kernel(Func kernelFunc) {
-  kernelFunc();
-}
-
 int main() {
   cl::sycl::accessor<F, 1, cl::sycl::access::mode::read_write> accessorA;
-    kernel<class kernel_function>(
+  cl::sycl::handler cgh;
+  cgh.single_task<class kernel_function>(
       [=]() {
         accessorA.use();
       });

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

@@ -0,0 +1,53 @@
+// RUN: %clang_cc1 -fsycl-is-device -verify -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.
+
+struct Base {
+  int X;
+};
+
+// This struct has non-standard layout, because both C (the most derived class)
+// and Base have non-static data members.
+struct C : public Base {
+  int Y;
+};
+
+template <typename name, typename Func>
+__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
+  kernelFunc();
+}
+
+
+void test() {
+  C C0;
+  C0.Y=0;
+  kernel_single_task<class MyKernel>([=] {
+    // expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
+    (void)C0.Y;
+  });
+}
+
+struct Kernel {
+  void operator()() {
+    (void) c1;
+    (void) c2;
+    (void) p;
+    (void) q;
+  }
+
+  int p;
+  // 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 'C'}}
+  C c2;
+};
+
+void test_struct_field() {
+  Kernel k{};
+
+  kernel_single_task<class kernel_object>(k);
+}

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

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only -std=c++11 %s
+// RUN: %clang_cc1 -fsycl-is-device -fsycl-new-kernel-param-requirements -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.

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

@@ -1,4 +1,4 @@
-// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only -std=c++11 %s
+// 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.