Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

[SYCL] Decompose kernel parameters and add inheritance support #1877

Merged
merged 60 commits into from
Jul 6, 2020
Merged
Changes from 1 commit
Commits
Show all changes
60 commits
Select commit Hold shift + click to select a range
7974c35
Support for arrays as kernel parameters.
rdeodhar Jun 9, 2020
4907194
Reusing some memberexpr building code.
rdeodhar Jun 9, 2020
d54c0ca
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 9, 2020
44d8663
Rebased changes from Elizabeth
Fznamznon Jun 10, 2020
801a0ea
Implement special bases handling
Fznamznon Jun 10, 2020
503638e
Merge branch 'sycl' into bases-handling
Fznamznon Jun 10, 2020
546c58d
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 10, 2020
52f2e5a
Handle vector types like scalar types
Fznamznon Jun 11, 2020
2a36a93
Fixed failing lit tests. Structs/Classes are no longer passed whole. …
elizabethandrews Jun 11, 2020
ab74fcf
Owner should be record being visited.
elizabethandrews Jun 12, 2020
4370d76
Avoid decomposing stream class. If field type is stream, we iterate
elizabethandrews Jun 12, 2020
52ce3f2
Updated support for arrays.
rdeodhar Jun 12, 2020
983b3d5
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 12, 2020
1bf0903
Formatting changes.
rdeodhar Jun 12, 2020
5d5121b
Formatting changes.
rdeodhar Jun 12, 2020
96ca8f4
Fix crash for stream type. Also changed handler call for consistency
elizabethandrews Jun 15, 2020
f03edd9
Correction to a test.
rdeodhar Jun 15, 2020
033b507
Merge remote-tracking branch 'rajiv_fork/akp2' into add_inheritance_s…
elizabethandrews Jun 15, 2020
4868d45
Fixed some crashes after merge.
elizabethandrews Jun 17, 2020
35383c5
Minor refactor
elizabethandrews Jun 17, 2020
de9e2aa
Added lit test for inheritance AST check
elizabethandrews Jun 17, 2020
d87b2cc
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 22, 2020
63cc362
Do not visit accessor fields
elizabethandrews Jun 23, 2020
51b598e
Do not generate default initializers for accessors in array
elizabethandrews Jun 23, 2020
2a1e9ba
Added CodeGen lit test
elizabethandrews Jun 25, 2020
0412db3
Array elements are now passed as individual parameters.
rdeodhar Jun 25, 2020
810af7b
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 25, 2020
48439c3
Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…
elizabethandrews Jun 25, 2020
d620e4e
Clang-Format Changes
elizabethandrews Jun 25, 2020
00c082f
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 25, 2020
af0b0c9
Corrections to temporarily disable tests expected to fail.
rdeodhar Jun 25, 2020
d5fb2d9
Changed tests to work with current array support.
rdeodhar Jun 26, 2020
70a2076
Fix multiple inheritance
Fznamznon Jun 25, 2020
f07c8d7
Add runtime test for functor inheritance
Fznamznon Jun 25, 2020
b1365c2
Add runtime test for accessor base
Fznamznon Jun 25, 2020
4ea6f47
Do not decompose cl::sycl::half type
Fznamznon Jun 25, 2020
15b47f4
Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…
elizabethandrews Jun 26, 2020
1c9e17b
Fix sampler lit test. Struct is decomposed.
elizabethandrews Jun 26, 2020
92e71bd
Cleaned up code a bit:
elizabethandrews Jun 26, 2020
f4cd574
Merge remote-tracking branch 'rajiv_fork/akp2' into add_inheritance_s…
elizabethandrews Jun 26, 2020
7bb1db5
ClangFormat Changes
elizabethandrews Jun 26, 2020
db492bd
Decomposed array elements, and changed manner of array element initia…
rdeodhar Jun 27, 2020
59cabac
Merge branch 'sycl' of https://github.com/otcshare/llvm into akp2
rdeodhar Jun 27, 2020
9f9b13d
Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…
elizabethandrews Jun 28, 2020
a64b209
Add CodeGen test for accessor inheritance
Fznamznon Jun 26, 2020
4c7dbd0
Add Sema AST test for accessor bases
Fznamznon Jun 26, 2020
1222a92
Merge remote-tracking branch 'rajiv_fork/akp2' into add_inheritance_s…
elizabethandrews Jun 29, 2020
6da194d
ClangFormat changes
elizabethandrews Jun 29, 2020
4afc3a3
Removed one redundant check.
rdeodhar Jun 29, 2020
1e5b360
Enable and fix array tests after merge.
elizabethandrews Jun 30, 2020
7af1020
Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…
elizabethandrews Jun 30, 2020
ed4d2f5
Merge remote-tracking branch 'rajiv_fork/akp2' into add_inheritance_s…
elizabethandrews Jun 30, 2020
4e1220a
Fix incorrect merge conflict resolution and ClangFormat error
elizabethandrews Jul 1, 2020
5dcf420
ClangFormat changes
elizabethandrews Jul 1, 2020
d5f56b3
Fix Windows test failure due to mangling
elizabethandrews Jul 2, 2020
7b81a3e
Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…
elizabethandrews Jul 2, 2020
91954fd
Removed unused variable and modified comments
elizabethandrews Jul 2, 2020
62ab84d
Fix incorrect merge resolution
elizabethandrews Jul 2, 2020
47d092a
Merge remote-tracking branch 'intel_llvm/sycl' into add_inheritance_s…
elizabethandrews Jul 2, 2020
a7ad39c
ClangFormat Change
elizabethandrews Jul 2, 2020
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
Rebased changes from Elizabeth
Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
  • Loading branch information
Fznamznon committed Jun 10, 2020
commit 44d866389576431f986ebe9a7b2ce7acef745f3a
134 changes: 92 additions & 42 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -683,9 +683,6 @@ constructKernelName(Sema &S, FunctionDecl *KernelCallerFunc,
// anonymous namespace so these don't get linkage.
namespace {

QualType getItemType(const FieldDecl *FD) { return FD->getType(); }
QualType getItemType(const CXXBaseSpecifier &BS) { return BS.getType(); }

// These enable handler execution only when previous handlers succeed.
template <typename... Tn>
static bool handleField(FieldDecl *FD, QualType FDTy, Tn &&... tn) {
Expand Down Expand Up @@ -729,11 +726,6 @@ template <typename T> using bind_param_t = typename bind_param<T>::type;
// })...)

// Implements the 'for-each-visitor' pattern.
template <typename ParentTy, typename... Handlers>
static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent,
CXXRecordDecl *Wrapper,
Handlers &... handlers);

template <typename RangeTy, typename... Handlers>
static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
Handlers &... handlers) {
Expand All @@ -742,7 +734,7 @@ static void VisitField(CXXRecordDecl *Owner, RangeTy &&Item, QualType ItemTy,
if (Util::isSyclStreamType(ItemTy))
KF_FOR_EACH(handleSyclStreamType, Item, ItemTy);
if (ItemTy->isStructureOrClassType())
VisitAccessorWrapper(Owner, Item, ItemTy->getAsCXXRecordDecl(),
VisitRecord(Owner, Item, ItemTy->getAsCXXRecordDecl(),
handlers...);
if (ItemTy->isArrayType())
VisitArrayElements(Item, ItemTy, handlers...);
Expand All @@ -762,38 +754,68 @@ static void VisitArrayElements(RangeTy Item, QualType FieldTy,
(void)std::initializer_list<int>{(handlers.leaveArray(ET, ElemCount), 0)...};
}

template <typename RangeTy, typename... Handlers>
static void VisitAccessorWrapperHelper(CXXRecordDecl *Owner, RangeTy Range,
Handlers &... handlers) {
for (const auto &Item : Range) {
QualType ItemTy = getItemType(Item);
(void)std::initializer_list<int>{(handlers.enterField(Owner, Item), 0)...};
VisitField(Owner, Item, ItemTy, handlers...);
(void)std::initializer_list<int>{(handlers.leaveField(Owner, Item), 0)...};
template <typename ParentTy, typename... Handlers>
static void VisitRecord(CXXRecordDecl *Owner, ParentTy &Parent,
CXXRecordDecl *Wrapper, Handlers &... handlers);

template <typename... Handlers>
static void VisitRecordHelper(CXXRecordDecl *Owner,
clang::CXXRecordDecl::base_class_range Range,
Handlers &... handlers) {
for (const auto &Base : Range) {
QualType BaseTy = Base.getType();
if (Util::isSyclAccessorType(BaseTy))
(void)std::initializer_list<int>{
(handlers.handleSyclAccessorType(Base, BaseTy), 0)...};
else if (Util::isSyclStreamType(BaseTy))
(void)std::initializer_list<int>{
(handlers.handleSyclStreamType(Base, BaseTy), 0)...};
else
VisitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), handlers...);
}
}

template <typename... Handlers>
static void VisitRecordHelper(CXXRecordDecl *Owner,
clang::RecordDecl::field_range Range,
Handlers &... handlers) {
VisitRecordFields(Owner, handlers...);
}

// Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter
// the Wrapper structure that we're currently visiting. Owner is the parent
// type (which doesn't exist in cases where it is a FieldDecl in the
// 'root'), and Wrapper is the current struct being unwrapped.
template <typename ParentTy, typename... Handlers>
static void VisitAccessorWrapper(CXXRecordDecl *Owner, ParentTy &Parent,
CXXRecordDecl *Wrapper,
Handlers &... handlers) {
static void VisitRecord(CXXRecordDecl *Owner, ParentTy &Parent,
CXXRecordDecl *Wrapper, Handlers &... handlers) {
(void)std::initializer_list<int>{(handlers.enterStruct(Owner, Parent), 0)...};
VisitAccessorWrapperHelper(Wrapper, Wrapper->bases(), handlers...);
VisitAccessorWrapperHelper(Wrapper, Wrapper->fields(), handlers...);
VisitRecordHelper(Wrapper, Wrapper->bases(), handlers...);
VisitRecordHelper(Wrapper, Wrapper->fields(), handlers...);
(void)std::initializer_list<int>{(handlers.leaveStruct(Owner, Parent), 0)...};
}

int getFieldNumber(const CXXRecordDecl *BaseDecl) {
int Members = 0;
for (const auto *Field : BaseDecl->fields())
++Members;

return Members;
}
elizabethandrews marked this conversation as resolved.
Show resolved Hide resolved

template <typename... Handlers>
static void VisitFunctorBases(CXXRecordDecl *KernelFunctor,
Handlers &... handlers) {
VisitRecordHelper(KernelFunctor, KernelFunctor->bases(), handlers...);
}


// A visitor function that dispatches to functions as defined in
// SyclKernelFieldHandler for the purposes of kernel generation.
template <typename... Handlers>
static void VisitRecordFields(RecordDecl::field_range Fields,
Handlers &... handlers) {
static void VisitRecordFields(CXXRecordDecl *Owner, Handlers &... handlers) {

for (const auto Field : Fields) {
for (const auto Field : Owner->fields()) {
(void)std::initializer_list<int>{
(handlers.enterField(nullptr, Field), 0)...};
QualType FieldTy = Field->getType();
Expand All @@ -807,12 +829,12 @@ static void VisitRecordFields(RecordDecl::field_range Fields,
else if (Util::isSyclStreamType(FieldTy)) {
// Stream actually wraps accessors, so do recursion
CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
VisitAccessorWrapper(nullptr, Field, RD, handlers...);
VisitRecord(nullptr, Field, RD, handlers...);
KF_FOR_EACH(handleSyclStreamType, Field, FieldTy);
} else if (FieldTy->isStructureOrClassType()) {
if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
VisitAccessorWrapper(nullptr, Field, RD, handlers...);
VisitRecord(nullptr, Field, RD, handlers...);
}
} else if (FieldTy->isReferenceType())
KF_FOR_EACH(handleReferenceType, Field, FieldTy);
Expand Down Expand Up @@ -1131,7 +1153,7 @@ class SyclKernelDeclCreator
}

bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy);
// addParam(FD, FieldTy);
return true;
}

Expand Down Expand Up @@ -1313,16 +1335,14 @@ class SyclKernelBodyCreator
bool handleSpecialType(FieldDecl *FD, QualType Ty) {
const auto *RecordDecl = Ty->getAsCXXRecordDecl();
// Perform initialization only if it is field of kernel object
if (MemberExprBases.size() == 1) {
InitializedEntity Entity =
InitializedEntity::InitializeMember(FD, &VarEntity);
// Initialize with the default constructor.
InitializationKind InitKind =
InitializationKind::CreateDefault(SourceLocation());
InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None);
InitExprs.push_back(MemberInit.get());
}
InitializedEntity Entity =
InitializedEntity::InitializeMember(FD, &VarEntity);
// Initialize with the default constructor.
InitializationKind InitKind =
InitializationKind::CreateDefault(SourceLocation());
InitializationSequence InitSeq(SemaRef, Entity, InitKind, None);
ExprResult MemberInit = InitSeq.Perform(SemaRef, Entity, InitKind, None);
InitExprs.push_back(MemberInit.get());
createSpecialMethodCall(RecordDecl, MemberExprBases.back(), InitMethodName,
FD);
return true;
Expand Down Expand Up @@ -1390,11 +1410,12 @@ class SyclKernelBodyCreator
}

bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
createExprForStructOrScalar(FD);
// createExprForStructOrScalar(FD);
return true;
}

bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
FieldTy->dump();
createExprForStructOrScalar(FD);
return true;
}
Expand All @@ -1403,8 +1424,35 @@ class SyclKernelBodyCreator
MemberExprBases.push_back(BuildMemberExpr(MemberExprBases.back(), FD));
}

void addStructInit(const CXXRecordDecl *RD){
if (!RD)
return;

int NumberOfFields = getFieldNumber(RD);
int popOut = NumberOfFields + RD->getNumBases();
llvm::SmallVector<Expr *, 16> BaseInitExprs;
for (int I = 0; I < popOut; I++) {
BaseInitExprs.push_back(InitExprs.back());
InitExprs.pop_back();
}
std::reverse(BaseInitExprs.begin(), BaseInitExprs.end());

Expr *ILE = new (SemaRef.getASTContext())
InitListExpr(SemaRef.getASTContext(), SourceLocation(), BaseInitExprs,
SourceLocation());
ILE->setType(QualType(RD->getTypeForDecl(), 0));
InitExprs.push_back(ILE);

}

void leaveStruct(const CXXRecordDecl *, FieldDecl *FD) final {
MemberExprBases.pop_back();
const CXXRecordDecl *RD = FD->getType()->getAsCXXRecordDecl();
addStructInit(RD);
}

void leaveStruct(const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) final {
const CXXRecordDecl *BaseClass = BS.getType()->getAsCXXRecordDecl();
addStructInit(BaseClass);
}

using SyclKernelFieldHandler::enterStruct;
Expand Down Expand Up @@ -1512,7 +1560,7 @@ class SyclKernelIntHeaderCreator
return true;
}
bool handleStructType(FieldDecl *FD, QualType FieldTy) final {
addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
// addParam(FD, FieldTy, SYCLIntegrationHeader::kind_std_layout);
return true;
}
bool handleScalarType(FieldDecl *FD, QualType FieldTy) final {
Expand Down Expand Up @@ -1606,7 +1654,9 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc,
StableName);

ConstructingOpenCLKernel = true;
VisitRecordFields(KernelLambda->fields(), checker, kernel_decl, kernel_body,
VisitFunctorBases(KernelLambda, checker, kernel_decl, kernel_body,
int_header);
VisitRecordFields(KernelLambda, checker, kernel_decl, kernel_body,
int_header);
ConstructingOpenCLKernel = false;
}
Expand Down