Skip to content

Commit 3adb4a5

Browse files
Fznamznonbader
Fznamznon
authored and
bader
committed
[SYCL] Adjust kernel parameters requirements (#1014)
Standard layout is too restrictive (prevents things like tuple). Enable standard layout requirement only if -sycl-std=1.2.1 is set. To make using of non-standard layout structs safe compiler builds structs layout in accordance with host ABI. Add trivially copyable requirement for all parameters. Signed-off-by: Mariya Podchishchaeva [email protected]
1 parent ed115ba commit 3adb4a5

File tree

13 files changed

+175
-55
lines changed

13 files changed

+175
-55
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -10314,8 +10314,10 @@ def err_sycl_virtual_types : Error<
1031410314
"No class with a vtable can be used in a SYCL kernel or any code included in the kernel">;
1031510315
def note_sycl_used_here : Note<"used here">;
1031610316
def note_sycl_recursive_function_declared_here: Note<"function implemented using recursion declared here">;
10317+
def err_sycl_non_trivially_copyable_type : Error<
10318+
"kernel parameter has non-trivially copyable class/struct type %0">;
1031710319
def err_sycl_non_std_layout_type : Error<
10318-
"kernel parameter has non-standard layout class/struct type">;
10320+
"kernel parameter has non-standard layout class/struct type %0">;
1031910321
def err_conflicting_sycl_kernel_attributes : Error<
1032010322
"conflicting attributes applied to a SYCL kernel">;
1032110323
def err_conflicting_sycl_function_attributes : Error<

‎clang/include/clang/Basic/LangOptions.def

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -233,6 +233,7 @@ LANGOPT(GPUMaxThreadsPerBlock, 32, 256, "default max threads per block for kerne
233233
LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device")
234234
LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation")
235235
LANGOPT(SYCLAllowFuncPtr , 1, 0, "Allow function pointers in SYCL device code")
236+
LANGOPT(SYCLStdLayoutKernelParams, 1, 0, "Enable standard layout requirement for SYCL kernel parameters")
236237
LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels")
237238

238239
LANGOPT(HIPUseNewLaunchAPI, 1, 0, "Use new kernel launching API for HIP")

‎clang/include/clang/Driver/CC1Options.td

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -899,6 +899,8 @@ def fsycl_int_header : Separate<["-"], "fsycl-int-header">,
899899
HelpText<"Generate SYCL integration header into this file.">;
900900
def fsycl_int_header_EQ : Joined<["-"], "fsycl-int-header=">,
901901
Alias<fsycl_int_header>;
902+
def fsycl_std_layout_kernel_params: Flag<["-"], "fsycl-std-layout-kernel-params">,
903+
HelpText<"Enable standard layout requirement for SYCL kernel parameters.">;
902904
def fsycl_allow_func_ptr : Flag<["-"], "fsycl-allow-func-ptr">,
903905
HelpText<"Allow function pointers in SYCL device.">;
904906
def fno_sycl_allow_func_ptr : Flag<["-"], "fno-sycl-allow-func-ptr">;

‎clang/lib/AST/RecordLayoutBuilder.cpp

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -2199,8 +2199,9 @@ static bool mustSkipTailPadding(TargetCXXABI ABI, const CXXRecordDecl *RD) {
21992199
llvm_unreachable("bad tail-padding use kind");
22002200
}
22012201

2202-
static bool isMsLayout(const ASTContext &Context) {
2203-
return Context.getTargetInfo().getCXXABI().isMicrosoft();
2202+
static bool isMsLayout(const ASTContext &Context, bool CheckAuxABI = false) {
2203+
return (CheckAuxABI) ? Context.getAuxTargetInfo()->getCXXABI().isMicrosoft()
2204+
: Context.getTargetInfo().getCXXABI().isMicrosoft();
22042205
}
22052206

22062207
// This section contains an implementation of struct layout that is, up to the
@@ -3025,6 +3026,9 @@ ASTContext::getASTRecordLayout(const RecordDecl *D) const {
30253026
// as soon as we begin to parse the definition. That definition is
30263027
// not a complete definition (which is what isDefinition() tests)
30273028
// until we *finish* parsing the definition.
3029+
bool CheckAuxABI = false;
3030+
if (getLangOpts().SYCLIsDevice && (getAuxTargetInfo() != nullptr))
3031+
CheckAuxABI = true;
30283032

30293033
if (D->hasExternalLexicalStorage() && !D->getDefinition())
30303034
getExternalSource()->CompleteType(const_cast<RecordDecl*>(D));
@@ -3042,7 +3046,7 @@ ASTContext::getASTRecordLayout(const RecordDecl *D) const {
30423046

30433047
const ASTRecordLayout *NewEntry = nullptr;
30443048

3045-
if (isMsLayout(*this)) {
3049+
if (isMsLayout(*this, CheckAuxABI)) {
30463050
MicrosoftRecordLayoutBuilder Builder(*this);
30473051
if (const auto *RD = dyn_cast<CXXRecordDecl>(D)) {
30483052
Builder.cxxLayout(RD);

‎clang/lib/Driver/ToolChains/Clang.cpp

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3989,6 +3989,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA,
39893989

39903990
if (Arg *A = Args.getLastArg(options::OPT_sycl_std_EQ)) {
39913991
A->render(Args, CmdArgs);
3992+
CmdArgs.push_back("-fsycl-std-layout-kernel-params");
39923993
} else if (IsSYCL) {
39933994
// Ensure the default version in SYCL mode is 1.2.1
39943995
CmdArgs.push_back("-sycl-std=1.2.1");

‎clang/lib/Frontend/CompilerInvocation.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -3130,6 +3130,8 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK,
31303130
Opts.SYCLIsHost = Args.hasArg(options::OPT_fsycl_is_host);
31313131
Opts.SYCLAllowFuncPtr = Args.hasFlag(options::OPT_fsycl_allow_func_ptr,
31323132
options::OPT_fno_sycl_allow_func_ptr, false);
3133+
Opts.SYCLStdLayoutKernelParams =
3134+
Args.hasArg(options::OPT_fsycl_std_layout_kernel_params);
31333135
Opts.SYCLUnnamedLambda = Args.hasArg(options::OPT_fsycl_unnamed_lambda);
31343136

31353137
// Set CUDA mode for OpenMP target NVPTX if specified in options

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 23 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -972,7 +972,6 @@ static target getAccessTarget(const ClassTemplateSpecializationDecl *AccTy) {
972972
// Returns true if all arguments are successfully built.
973973
static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
974974
SmallVectorImpl<ParamDesc> &ParamDescs) {
975-
const LambdaCapture *Cpt = KernelObj->captures_begin();
976975
auto CreateAndAddPrmDsc = [&](const FieldDecl *Fld, const QualType &ArgType) {
977976
// Create a parameter descriptor and append it to the result
978977
ParamDescs.push_back(makeParamDesc(Fld, ArgType));
@@ -1036,26 +1035,35 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
10361035
QualType ArgTy = Fld->getType();
10371036
if (Util::isSyclAccessorType(ArgTy) || Util::isSyclSamplerType(ArgTy)) {
10381037
createSpecialSYCLObjParamDesc(Fld, ArgTy);
1039-
} else if (!ArgTy->isStandardLayoutType()) {
1040-
// SYCL v1.2.1 s4.8.10 p5:
1041-
// C++ non-standard layout values must not be passed as arguments to a
1042-
// kernel that is compiled for a device.
1043-
const auto &DiagLocation =
1044-
Cpt ? Cpt->getLocation() : cast<DeclaratorDecl>(Fld)->getLocation();
1045-
1046-
Context.getDiagnostics().Report(DiagLocation,
1047-
diag::err_sycl_non_std_layout_type);
1048-
1049-
// Set the flag and continue processing so we can emit error for each
1050-
// invalid argument.
1051-
AllArgsAreValid = false;
10521038
} else if (ArgTy->isStructureOrClassType()) {
1053-
assert(ArgTy->isStandardLayoutType());
1039+
if (Context.getLangOpts().SYCLStdLayoutKernelParams) {
1040+
if (!ArgTy->isStandardLayoutType()) {
1041+
Context.getDiagnostics().Report(Fld->getLocation(),
1042+
diag::err_sycl_non_std_layout_type)
1043+
<< ArgTy;
1044+
AllArgsAreValid = false;
1045+
continue;
1046+
}
1047+
}
1048+
// TODO: Make stream class trivially copyable and remove the check on
1049+
// stream class.
1050+
if (!ArgTy.isTriviallyCopyableType(Context) &&
1051+
!Util::isSyclStreamType(ArgTy)) {
1052+
Context.getDiagnostics().Report(
1053+
Fld->getLocation(), diag::err_sycl_non_trivially_copyable_type)
1054+
<< ArgTy;
1055+
AllArgsAreValid = false;
1056+
continue;
1057+
}
10541058

10551059
CreateAndAddPrmDsc(Fld, ArgTy);
10561060

10571061
// Create descriptors for each accessor field in the class or struct
10581062
createParamDescForWrappedAccessors(Fld, ArgTy);
1063+
} else if (ArgTy->isReferenceType()) {
1064+
Context.getDiagnostics().Report(
1065+
Fld->getLocation(), diag::err_bad_kernel_param_type) << ArgTy;
1066+
AllArgsAreValid = false;
10591067
} else if (ArgTy->isPointerType()) {
10601068
// Pointer Arguments need to be in the global address space
10611069
QualType PointeeTy = ArgTy->getPointeeType();
@@ -1071,10 +1079,6 @@ static bool buildArgTys(ASTContext &Context, CXXRecordDecl *KernelObj,
10711079
} else {
10721080
llvm_unreachable("Unsupported kernel parameter type");
10731081
}
1074-
1075-
// Update capture iterator as we process arguments
1076-
if (Cpt && Cpt != KernelObj->captures_end())
1077-
++Cpt;
10781082
}
10791083

10801084
return AllArgsAreValid;

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

Lines changed: 21 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,21 @@
1+
// 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
2+
// 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
3+
4+
#include "sycl.hpp"
5+
// CHK-WIN: %struct{{.*}}F = type { i8, i8 }
6+
// CHK-LIN: %struct{{.*}}F = type { i8 }
7+
struct F1 {};
8+
struct F2 {};
9+
struct F : F1, F2 {
10+
char x;
11+
};
12+
13+
int main() {
14+
cl::sycl::accessor<F, 1, cl::sycl::access::mode::read_write> accessorA;
15+
cl::sycl::handler cgh;
16+
cgh.single_task<class kernel_function>(
17+
[=]() {
18+
accessorA.use();
19+
});
20+
return 0;
21+
}

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

Lines changed: 5 additions & 31 deletions
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,5 @@
1-
// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
1+
// RUN: %clang_cc1 -fsycl-is-device -fsycl-std-layout-kernel-params -verify -fsyntax-only %s
2+
// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only %s
23

34
// This test checks if compiler reports compilation error on an attempt to pass
45
// non-standard layout struct object as SYCL kernel parameter.
@@ -23,38 +24,11 @@ void test() {
2324
C C0;
2425
C0.Y=0;
2526
kernel_single_task<class MyKernel>([=] {
26-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
27+
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
2728
(void)C0.Y;
2829
});
2930
}
3031

31-
void test_capture_explicit_ref() {
32-
int p = 0;
33-
double q = 0;
34-
float s = 0;
35-
kernel_single_task<class kernel_capture_single_ref>([
36-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
37-
&p,
38-
q,
39-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
40-
&s] {
41-
(void) q;
42-
(void) p;
43-
(void) s;
44-
});
45-
}
46-
47-
void test_capture_implicit_refs() {
48-
int p = 0;
49-
double q = 0;
50-
kernel_single_task<class kernel_capture_refs>([&] {
51-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
52-
(void) p;
53-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
54-
(void) q;
55-
});
56-
}
57-
5832
struct Kernel {
5933
void operator()() {
6034
(void) c1;
@@ -64,12 +38,12 @@ struct Kernel {
6438
}
6539

6640
int p;
67-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
41+
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
6842
C c1;
6943

7044
int q;
7145

72-
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type}}
46+
// expected-error@+1 {{kernel parameter has non-standard layout class/struct type 'C'}}
7347
C c2;
7448
};
7549

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

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,28 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
2+
3+
// This test checks if compiler reports compilation error on an attempt to pass
4+
// a struct with non-trivially copyable type as SYCL kernel parameter.
5+
6+
struct A { int i; };
7+
8+
struct B {
9+
int i;
10+
B (int _i) : i(_i) {}
11+
B (const B& x) : i(x.i) {}
12+
};
13+
14+
template <typename Name, typename Func>
15+
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
16+
kernelFunc();
17+
}
18+
19+
void test() {
20+
A IamGood;
21+
IamGood.i = 0;
22+
B IamBad(1);
23+
kernel_single_task<class kernel_capture_refs>([=] {
24+
int a = IamGood.i;
25+
// expected-error@+1 {{kernel parameter has non-trivially copyable class/struct type}}
26+
int b = IamBad.i;
27+
});
28+
}

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

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,36 @@
1+
// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s
2+
3+
// This test checks if compiler reports compilation error on an attempt to pass
4+
// a reference as SYCL kernel parameter.
5+
6+
template <typename Name, typename Func>
7+
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
8+
kernelFunc();
9+
}
10+
11+
void test_capture_explicit_ref() {
12+
int p = 0;
13+
double q = 0;
14+
float s = 0;
15+
kernel_single_task<class kernel_capture_single_ref>([
16+
// expected-error@+1 {{'int &' cannot be used as the type of a kernel parameter}}
17+
&p,
18+
q,
19+
// expected-error@+1 {{'float &' cannot be used as the type of a kernel parameter}}
20+
&s] {
21+
(void) q;
22+
(void) p;
23+
(void) s;
24+
});
25+
}
26+
27+
void test_capture_implicit_refs() {
28+
int p = 0;
29+
double q = 0;
30+
kernel_single_task<class kernel_capture_refs>([&] {
31+
// expected-error@+1 {{'int &' cannot be used as the type of a kernel parameter}}
32+
(void) p;
33+
// expected-error@+1 {{'double &' cannot be used as the type of a kernel parameter}}
34+
(void) q;
35+
});
36+
}

‎sycl/test/basic_tests/handler/handler_mem_op.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -19,7 +19,7 @@
1919
using namespace cl::sycl;
2020

2121
template <typename T> struct point {
22-
point(const point &rhs) : x(rhs.x), y(rhs.y) {}
22+
point(const point &rhs) = default;
2323
point(T x, T y) : x(x), y(y) {}
2424
point(T v) : x(v), y(v) {}
2525
point() : x(0), y(0) {}

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

Lines changed: 45 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,45 @@
1+
// RUN: %clangxx -fsycl %s -o %t.out
2+
// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
3+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
4+
// RUN: %GPU_RUN_PLACEHOLDER %t.out
5+
// RUN: %ACC_RUN_PLACEHOLDER %t.out
6+
7+
#include <CL/sycl.hpp>
8+
9+
using namespace cl::sycl;
10+
11+
struct F1 {};
12+
struct F2 {};
13+
struct F : F1, F2 {
14+
cl::sycl::cl_char x;
15+
};
16+
17+
bool test0() {
18+
F S;
19+
S.x = 0;
20+
F S0;
21+
S0.x = 1;
22+
{
23+
buffer<F, 1> Buf(&S0, range<1>(1));
24+
queue myQueue;
25+
myQueue.submit([&](handler &cgh) {
26+
auto B = Buf.get_access<access::mode::write>(cgh);
27+
cgh.single_task<class NonStandard>([=] { B[0] = S; });
28+
});
29+
}
30+
bool Passed = (S.x == S0.x);
31+
32+
if (!Passed) {
33+
std::cout << "test0 failed" << std::endl;
34+
}
35+
return Passed;
36+
}
37+
38+
int main() {
39+
40+
bool Pass = test0();
41+
42+
std::cout << "Test " << (Pass ? "passed" : "FAILED") << std::endl;
43+
return Pass ? 0 : 1;
44+
45+
}

0 commit comments

Comments
 (0)