Skip to content

Commit b8b9c2c

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 6434b9c + 038d002 commit b8b9c2c

37 files changed

+885
-573
lines changed

‎clang/include/clang/Basic/Attr.td

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1246,16 +1246,16 @@ def SYCLIntelNumSimdWorkItems : InheritableAttr {
12461246
let PragmaAttributeSupport = 0;
12471247
}
12481248

1249-
def SYCLIntelStallEnable : InheritableAttr {
1250-
let Spellings = [CXX11<"intel","stall_enable">];
1249+
def SYCLIntelUseStallEnableClusters : InheritableAttr {
1250+
let Spellings = [CXX11<"intel","use_stall_enable_clusters">];
12511251
let LangOpts = [SYCLIsHost, SYCLIsDevice];
12521252
let Subjects = SubjectList<[Function], ErrorDiag>;
12531253
let AdditionalMembers = [{
12541254
static const char *getName() {
12551255
return "stall_enable";
12561256
}
12571257
}];
1258-
let Documentation = [SYCLIntelStallEnableAttrDocs];
1258+
let Documentation = [SYCLIntelUseStallEnableClustersAttrDocs];
12591259
let PragmaAttributeSupport = 0;
12601260
}
12611261

‎clang/include/clang/Basic/AttrDocs.td

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -2241,33 +2241,33 @@ device kernel, the attribute is ignored and it is not propagated to a kernel.
22412241
}];
22422242
}
22432243

2244-
def SYCLIntelStallEnableAttrDocs : Documentation {
2244+
def SYCLIntelUseStallEnableClustersAttrDocs : Documentation {
22452245
let Category = DocCatFunction;
2246-
let Heading = "intel::stall_enable";
2246+
let Heading = "intel::use_stall_enable_clusters";
22472247
let Content = [{
22482248
When applied to a lambda or function call operator (of a function object)
22492249
on device, this requests, to the extent possible, that statically-scheduled
22502250
clusters handle stalls using a stall-enable signal to freeze computation
22512251
within the cluster. This attribute is ignored on the host.
22522252

2253-
If ``intel::stall_enable`` is applied to a function called from a device
2253+
If ``intel::use_stall_enable_clusters`` is applied to a function called from a device
22542254
kernel, the attribute is ignored and it is not propagated to a kernel.
22552255

2256-
The ``intel::stall_enable`` attribute takes no argument and has an effect
2256+
The ``intel::use_stall_enable_clusters`` attribute takes no argument and has an effect
22572257
when applied to a function, and no effect otherwise.
22582258

22592259
.. code-block:: c++
22602260

22612261
class Functor
22622262
{
2263-
[[intel::stall_enable]] void operator()(item<1> item)
2263+
[[intel::use_stall_enable_clusters]] void operator()(item<1> item)
22642264
{
22652265
/* kernel code */
22662266
}
22672267
}
22682268

22692269
kernel<class kernel_name>(
2270-
[]() [[intel::stall_enable]] {
2270+
[]() [[intel::use_stall_enable_clusters]] {
22712271
/* kernel code */
22722272
});
22732273

‎clang/include/clang/Basic/AttributeCommonInfo.h

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -166,7 +166,7 @@ class AttributeCommonInfo {
166166
ParsedAttr == AT_SYCLIntelMaxWorkGroupSize ||
167167
ParsedAttr == AT_SYCLIntelMaxGlobalWorkDim ||
168168
ParsedAttr == AT_SYCLIntelNoGlobalWorkOffset ||
169-
ParsedAttr == AT_SYCLIntelStallEnable)
169+
ParsedAttr == AT_SYCLIntelUseStallEnableClusters)
170170
return true;
171171

172172
return false;

‎clang/lib/CodeGen/CodeGenFunction.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -694,7 +694,7 @@ void CodeGenFunction::EmitOpenCLKernelMetadata(const FunctionDecl *FD,
694694
Fn->setMetadata("no_global_work_offset", llvm::MDNode::get(Context, {}));
695695
}
696696

697-
if (FD->hasAttr<SYCLIntelStallEnableAttr>()) {
697+
if (FD->hasAttr<SYCLIntelUseStallEnableClustersAttr>()) {
698698
llvm::Metadata *AttrMDArgs[] = {
699699
llvm::ConstantAsMetadata::get(Builder.getInt32(1))};
700700
Fn->setMetadata("stall_enable", llvm::MDNode::get(Context, AttrMDArgs));

‎clang/lib/Sema/SemaDeclAttr.cpp

Lines changed: 6 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -3030,8 +3030,9 @@ static void handleNumSimdWorkItemsAttr(Sema &S, Decl *D,
30303030
E);
30313031
}
30323032

3033-
// Handles stall_enable
3034-
static void handleStallEnableAttr(Sema &S, Decl *D, const ParsedAttr &Attr) {
3033+
// Handles use_stall_enable_clusters
3034+
static void handleUseStallEnableClustersAttr(Sema &S, Decl *D,
3035+
const ParsedAttr &Attr) {
30353036
if (D->isInvalidDecl())
30363037
return;
30373038

@@ -3041,7 +3042,7 @@ static void handleStallEnableAttr(Sema &S, Decl *D, const ParsedAttr &Attr) {
30413042
return;
30423043
}
30433044

3044-
handleSimpleAttribute<SYCLIntelStallEnableAttr>(S, D, Attr);
3045+
handleSimpleAttribute<SYCLIntelUseStallEnableClustersAttr>(S, D, Attr);
30453046
}
30463047

30473048
// Add scheduler_target_fmax_mhz
@@ -8417,8 +8418,8 @@ static void ProcessDeclAttribute(Sema &S, Scope *scope, Decl *D,
84178418
case ParsedAttr::AT_SYCLIntelNoGlobalWorkOffset:
84188419
handleNoGlobalWorkOffsetAttr(S, D, AL);
84198420
break;
8420-
case ParsedAttr::AT_SYCLIntelStallEnable:
8421-
handleStallEnableAttr(S, D, AL);
8421+
case ParsedAttr::AT_SYCLIntelUseStallEnableClusters:
8422+
handleUseStallEnableClustersAttr(S, D, AL);
84228423
break;
84238424
case ParsedAttr::AT_VecTypeHint:
84248425
handleVecTypeHint(S, D, AL);

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 10 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -547,16 +547,16 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
547547
if (auto *A = FD->getAttr<SYCLSimdAttr>())
548548
Attrs.insert(A);
549549

550-
// Allow the kernel attribute "stall_enable" only on lambda functions
551-
// and function objects that are called directly from a kernel
550+
// Allow the kernel attribute "use_stall_enable_clusters" only on lambda
551+
// functions and function objects that are called directly from a kernel
552552
// (i.e. the one passed to the single_task or parallel_for functions).
553553
// For all other cases, emit a warning and ignore.
554-
if (auto *A = FD->getAttr<SYCLIntelStallEnableAttr>()) {
554+
if (auto *A = FD->getAttr<SYCLIntelUseStallEnableClustersAttr>()) {
555555
if (ParentFD == SYCLKernel) {
556556
Attrs.insert(A);
557557
} else {
558558
SemaRef.Diag(A->getLocation(), diag::warn_attribute_ignored) << A;
559-
FD->dropAttr<SYCLIntelStallEnableAttr>();
559+
FD->dropAttr<SYCLIntelUseStallEnableClustersAttr>();
560560
}
561561
}
562562

@@ -1214,7 +1214,9 @@ void KernelObjVisitor::visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
12141214
const CXXRecordDecl *Wrapper,
12151215
QualType RecordTy,
12161216
HandlerTys &... Handlers) {
1217-
if (RecordTy->getAsRecordDecl()->hasAttr<SYCLRequiresDecompositionAttr>()) {
1217+
RecordDecl *RD = RecordTy->getAsRecordDecl();
1218+
assert(RD && "should not be null.");
1219+
if (RD->hasAttr<SYCLRequiresDecompositionAttr>()) {
12181220
// If this container requires decomposition, we have to visit it as
12191221
// 'complex', so all handlers are called in this case with the 'complex'
12201222
// case.
@@ -1583,6 +1585,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
15831585
bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final {
15841586
if (CollectionStack.pop_back_val()) {
15851587
RecordDecl *RD = Ty->getAsRecordDecl();
1588+
assert(RD && "should not be null.");
15861589
if (!RD->hasAttr<SYCLRequiresDecompositionAttr>())
15871590
RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit(
15881591
SemaRef.getASTContext()));
@@ -1601,6 +1604,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler {
16011604
QualType Ty) final {
16021605
if (CollectionStack.pop_back_val()) {
16031606
RecordDecl *RD = Ty->getAsRecordDecl();
1607+
assert(RD && "should not be null.");
16041608
if (!RD->hasAttr<SYCLRequiresDecompositionAttr>())
16051609
RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit(
16061610
SemaRef.getASTContext()));
@@ -3277,7 +3281,7 @@ void Sema::MarkDevice(void) {
32773281
case attr::Kind::SYCLIntelSchedulerTargetFmaxMhz:
32783282
case attr::Kind::SYCLIntelMaxGlobalWorkDim:
32793283
case attr::Kind::SYCLIntelNoGlobalWorkOffset:
3280-
case attr::Kind::SYCLIntelStallEnable:
3284+
case attr::Kind::SYCLIntelUseStallEnableClusters:
32813285
case attr::Kind::SYCLSimd: {
32823286
if ((A->getKind() == attr::Kind::SYCLSimd) && KernelBody &&
32833287
!KernelBody->getAttr<SYCLSimdAttr>()) {

‎clang/test/CodeGenSYCL/stall_enable.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -7,7 +7,7 @@ queue q;
77

88
class Foo {
99
public:
10-
[[intel::stall_enable]] void operator()() const {}
10+
[[intel::use_stall_enable_clusters]] void operator()() const {}
1111
};
1212

1313
int main() {
@@ -16,7 +16,7 @@ int main() {
1616
h.single_task<class test_kernel1>(f);
1717

1818
h.single_task<class test_kernel2>(
19-
[]() [[intel::stall_enable]]{});
19+
[]() [[intel::use_stall_enable_clusters]]{});
2020
});
2121
return 0;
2222
}

‎clang/test/SemaSYCL/accessor_inheritance.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -42,8 +42,8 @@ int main() {
4242
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_A' 'int'
4343
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
4444
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_B' 'int'
45-
// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>' 'void () noexcept'
46-
// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>' 'void () noexcept'
45+
// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
46+
// CHECK-NEXT: CXXConstructExpr {{.*}} 'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read>':'cl::sycl::accessor<char, 1, cl::sycl::access::mode::read, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
4747
// CHECK-NEXT: ImplicitCastExpr {{.*}} 'int' <LValueToRValue>
4848
// CHECK-NEXT: DeclRefExpr {{.*}} 'int' lvalue ParmVar {{.*}} '_arg_C' 'int'
4949

‎clang/test/SemaSYCL/basic-kernel-wrapper.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ int main() {
4343

4444
// CHECK: CXXMemberCallExpr {{.*}} 'void'
4545
// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
46-
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>' lvalue .
46+
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' lvalue .
4747
// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}basic-kernel-wrapper.cpp{{.*}})' lvalue Var
4848

4949
// CHECK-NEXT: ImplicitCastExpr {{.*}} <LValueToRValue>

‎clang/test/SemaSYCL/implicit_kernel_type.cpp

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,6 +1,6 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -DERROR
2-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsyntax-only -sycl-std=2020 -verify %s -Wsycl-strict -DWARN
3-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict -DERROR
2+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify %s -Wsycl-strict -DWARN
3+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -internal-isystem %S/Inputs -fsycl-unnamed-lambda -fsyntax-only -sycl-std=2020 -verify %s -Werror=sycl-strict
44

55
#include "sycl.hpp"
66

‎clang/test/SemaSYCL/kernelname-enum.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,4 +1,4 @@
1-
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s
1+
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s
22

33
#include "Inputs/sycl.hpp"
44

‎clang/test/SemaSYCL/stall_enable.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -6,31 +6,31 @@
66
using namespace cl::sycl;
77
queue q;
88

9-
[[intel::stall_enable]] void test() {} //expected-warning{{'stall_enable' attribute ignored}}
9+
[[intel::use_stall_enable_clusters]] void test() {} //expected-warning{{'use_stall_enable_clusters' attribute ignored}}
1010

1111
#ifdef TRIGGER_ERROR
12-
[[intel::stall_enable(1)]] void bar1() {} // expected-error{{'stall_enable' attribute takes no arguments}}
13-
[[intel::stall_enable]] int N; // expected-error{{'stall_enable' attribute only applies to functions}}
12+
[[intel::use_stall_enable_clusters(1)]] void bar1() {} // expected-error{{'use_stall_enable_clusters' attribute takes no arguments}}
13+
[[intel::use_stall_enable_clusters]] int N; // expected-error{{'use_stall_enable_clusters' attribute only applies to functions}}
1414
#endif
1515

1616
struct FuncObj {
17-
[[intel::stall_enable]] void operator()() const {}
17+
[[intel::use_stall_enable_clusters]] void operator()() const {}
1818
};
1919

2020
int main() {
2121
q.submit([&](handler &h) {
2222
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel1
23-
// CHECK: SYCLIntelStallEnableAttr {{.*}}
23+
// CHECK: SYCLIntelUseStallEnableClustersAttr {{.*}}
2424
h.single_task<class test_kernel1>(
2525
FuncObj());
2626

2727
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel2
28-
// CHECK: SYCLIntelStallEnableAttr {{.*}}
28+
// CHECK: SYCLIntelUseStallEnableClustersAttr {{.*}}
2929
h.single_task<class test_kernel2>(
30-
[]() [[intel::stall_enable]]{});
30+
[]() [[intel::use_stall_enable_clusters]]{});
3131

3232
// CHECK-LABEL: FunctionDecl {{.*}}test_kernel3
33-
// CHECK-NOT: SYCLIntelStallEnableAttr {{.*}}
33+
// CHECK-NOT: SYCLIntelUseStallEnableClustersAttr {{.*}}
3434
h.single_task<class test_kernel3>(
3535
[]() { test(); });
3636
});

‎clang/test/SemaSYCL/wrapped-accessor.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -35,14 +35,14 @@ int main() {
3535
// argument
3636
// CHECK: VarDecl {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
3737
// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
38-
// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>>'
39-
// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>' 'void () noexcept'
38+
// CHECK-NEXT: InitListExpr {{.*}}'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>'
39+
// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' 'void () noexcept'
4040

4141
// Check that accessor field of the wrapper object is initialized using __init method
4242
// CHECK-NEXT: CXXMemberCallExpr {{.*}} 'void'
4343
// CHECK-NEXT: MemberExpr {{.*}} 'void ({{.*}}PtrType, range<1>, range<1>, id<1>)' lvalue .__init
44-
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>' lvalue .accessor {{.*}}
45-
// CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t, cl::sycl::ONEAPI::accessor_property_list<>>>' lvalue .
44+
// CHECK-NEXT: MemberExpr {{.*}} 'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>':'cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>' lvalue .accessor {{.*}}
45+
// CHECK-NEXT: MemberExpr {{.*}} 'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>':'AccWrapper<cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write, cl::sycl::access::target::global_buffer, cl::sycl::access::placeholder::false_t>>' lvalue .
4646
// CHECK-NEXT: DeclRefExpr {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})' lvalue Var {{.*}} '(lambda at {{.*}}wrapped-accessor.cpp{{.*}})'
4747

4848
// Parameters of the _init method

‎sycl/doc/ABIPolicyGuide.md

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -87,10 +87,11 @@ There is a set of tests to help identifying ABI changes:
8787
Please, prefer updating the test files with the above command. The checker
8888
script automatically sorts symbols. This would allow developers to avoid
8989
large diffs and help maintainers identify the nature of ABI changes.
90-
* `test/abi/layout*` and `test/abi/symbol_size*` are a group of tests to check
91-
the internal layout of some classes. The `layout*` tests check some of API
92-
classes for layout changes, while `symbol_size` only checks `sizeof` for API
93-
classes. Changing the class layout is a breaking change.
90+
* `test/abi/layout*` and `test/abi/symbol_size_alignment.cpp` are a group of
91+
tests to check the internal layout of some classes. The `layout*` tests check
92+
some of API classes for layout changes, while `symbol_size_alignment` only
93+
checks `sizeof` and `alignof` for API classes. Changing the class layout is a
94+
breaking change.
9495

9596
## Changing ABI
9697

‎sycl/doc/CompilerAndRuntimeDesign.md

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -498,7 +498,9 @@ unit)
498498
TBD
499499

500500
##### Specialization constants lowering
501-
TBD
501+
502+
See [corresponding documentation](SpecializationConstants.md)
503+
502504

503505
#### CUDA support
504506

0 commit comments

Comments
 (0)