Skip to content

[SYCL] moving type checks to later in Semantic Analysis lifecycle #1465

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.

Sign up for GitHub

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

Merged
merged 15 commits into from
Apr 8, 2020
Merged

[SYCL] moving type checks to later in Semantic Analysis lifecycle #1465

Show file tree
Hide file tree
Changes from 7 commits
Commits
Show all changes
15 commits
Select commit Hold shift + click to select a range
c89f95d
rebased to get past merge conflict
cperkinsintel Apr 2, 2020
961636f
moving SYCL type checks out of SemaType so they are called later in t…
cperkinsintel Apr 2, 2020
4fc329d
clang format
cperkinsintel Apr 2, 2020
aadb01e
more clang format
cperkinsintel Apr 2, 2020
4aa0cf2
feedback from PR
cperkinsintel Apr 2, 2020
c568ed7
testing changes after resolving merge conflict and more PR feedback
cperkinsintel Apr 3, 2020
bb43284
clang format (no way)
cperkinsintel Apr 3, 2020
eae7e2d
incorporate feedback from PR. Added recursive struct member type chec…
cperkinsintel Apr 3, 2020
3b497d3
incorporating feedback from PR
cperkinsintel Apr 6, 2020
ff71ec7
flang-cormat
cperkinsintel Apr 6, 2020
9bef479
more small changes
cperkinsintel Apr 6, 2020
0d63fc9
merged typecase checsk
cperkinsintel Apr 7, 2020
a53dcfb
braces
cperkinsintel Apr 7, 2020
8e59a2c
tighten recursion code, add function proto test, update comment
cperkinsintel Apr 7, 2020
2a4ada4
recorDecl instead of recordCXXDecl
cperkinsintel Apr 7, 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
1 change: 1 addition & 0 deletions clang/include/clang/Sema/Sema.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12455,6 +12455,7 @@ class Sema final {
};

bool isKnownGoodSYCLDecl(const Decl *D);
void checkSYCLVarDeclIfInKernel(VarDecl *Var);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Still kind of hate the name, seems clumsy, but I don't have a better suggestion at the moment, I'll think about it.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suggested change
void checkSYCLVarDeclIfInKernel(VarDecl *Var);
void checkSYCLVarDeclInSYCLKernel(VarDecl *Var);

since SYCL is about the kernel, not the variable?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

SYCL is not only about kernels. Let's call it "device code".

Suggested change
void checkSYCLVarDeclIfInKernel(VarDecl *Var);
void diagnoseVarDeclIfSYCLDeviceCode(VarDecl *Var);

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ping?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

How about 'checkSYCLDeviceVarDecl'. "check" is the common word for a function that diagnoses. The word "if" is incorrect, since this isn't checking whether it is in a SYCL device.

The above also is more consistent in the cases we've added where we recursively check.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Hmm, okay. Let it be checkSYCLDeviceVarDecl.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

done.

void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC);
void MarkDevice();

Expand Down
3 changes: 3 additions & 0 deletions clang/lib/Sema/SemaDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12660,6 +12660,9 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) {
}
}

if (getLangOpts().SYCLIsDevice)
checkSYCLVarDeclIfInKernel(var);

// In Objective-C, don't allow jumps past the implicit initialization of a
// local retaining variable.
if (getLangOpts().ObjC &&
Expand Down
38 changes: 35 additions & 3 deletions clang/lib/Sema/SemaSYCL.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -200,6 +200,41 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) {
return false;
}

bool isZeroSizedArray(QualType Ty) {
if (const auto *CATy = dyn_cast<ConstantArrayType>(Ty)) {
return (CATy->getSize() == 0);
}
return false;
}

void Sema::checkSYCLVarDeclIfInKernel(VarDecl *Var) {
// not all variable types supported in kernel contexts
// if not we record a deferred diagnostic.
assert(getLangOpts().SYCLIsDevice &&
"Should only be called during SYCL compilation");
QualType Ty = Var->getType();
SourceRange Loc = Var->getLocation();

// __int128, __int128_t, __uint128_t
if (Ty->isSpecificBuiltinType(BuiltinType::Int128) ||
Ty->isSpecificBuiltinType(BuiltinType::UInt128))
SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported)
<< Ty.getUnqualifiedType().getCanonicalType().getAsString();

// QuadType __float128
if (Ty->isSpecificBuiltinType(BuiltinType::Float128) &&
!Context.getTargetInfo().hasFloat128Type())
SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported)
<< "__float128";

// zero length arrays
if (isZeroSizedArray(Ty))
SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size);

// TODO: check type of accessor
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

??

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

BTW I have a feeling that it is not a proper place for accessor type checking. Just because accessors always declared in host code and captured to the device code. We can instead:

  • Check type of buffer/accessor data using static assertions in SYCL headers.
  • Check type of accessors while we iterate over kernel object captures in SemaSYCL (where check for trivially destructible/constructible kernel arguments here).

I personally prefer the way with SYCL headers, since SemaSYCL extremely big, complicated and ugly.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ping?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'll remove the comment. This is the next thing I am planning to undertake, but not as part of this PR.

// if(Util::isSyclAccessorType(Ty))
}

class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
public:
MarkDeviceFunction(Sema &S)
Expand Down Expand Up @@ -229,7 +264,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
if (Method->isVirtual())
SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict)
<< Sema::KernelCallVirtualFunction;

CheckSYCLType(Callee->getReturnType(), Callee->getSourceRange());

if (auto const *FD = dyn_cast<FunctionDecl>(Callee)) {
Expand Down Expand Up @@ -300,7 +334,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
Decl *D = E->getDecl();
if (SemaRef.isKnownGoodSYCLDecl(D))
return true;

CheckSYCLType(E->getType(), E->getSourceRange());
return true;
}
Expand Down Expand Up @@ -439,7 +472,6 @@ class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
SemaRef.Diag(Loc.getBegin(), diag::err_vla_unsupported);
return false;
}

while (Ty->isAnyPointerType() || Ty->isArrayType())
Ty = QualType{Ty->getPointeeOrArrayElementType(), 0};

Expand Down
14 changes: 2 additions & 12 deletions clang/lib/Sema/SemaType.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1527,12 +1527,8 @@ static QualType ConvertDeclSpecToType(TypeProcessingState &state) {
break;
case DeclSpec::TST_float128:
if (!S.Context.getTargetInfo().hasFloat128Type() &&
S.getLangOpts().SYCLIsDevice)
S.SYCLDiagIfDeviceCode(DS.getTypeSpecTypeLoc(),
diag::err_type_unsupported)
<< "__float128";
else if (!S.Context.getTargetInfo().hasFloat128Type() &&
!(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
!S.getLangOpts().SYCLIsDevice &&
!(S.getLangOpts().OpenMP && S.getLangOpts().OpenMPIsDevice))
S.Diag(DS.getTypeSpecTypeLoc(), diag::err_type_unsupported)
<< "__float128";
Result = Context.Float128Ty;
Expand Down Expand Up @@ -2350,12 +2346,6 @@ QualType Sema::BuildArrayType(QualType T, ArrayType::ArraySizeModifier ASM,
<< ArraySize->getSourceRange();
ASM = ArrayType::Normal;
}

// Zero length arrays are disallowed in SYCL device code.
if (getLangOpts().SYCLIsDevice)
SYCLDiagIfDeviceCode(ArraySize->getBeginLoc(),
diag::err_typecheck_zero_array_size)
<< ArraySize->getSourceRange();
} else if (!T->isDependentType() && !T->isVariablyModifiedType() &&
!T->isIncompleteType() && !T->isUndeducedType()) {
// Is the array too large?
Expand Down
103 changes: 95 additions & 8 deletions clang/test/SemaSYCL/deferred-diagnostics-emit.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,13 +2,18 @@
//
// Ensure that the SYCL diagnostics that are typically deferred are correctly emitted.

namespace std {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do you validate variable templates anywhere? How about alias templates?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I put those over in sycl-restrict.cpp, along with checks for auto, typedef, and some false postives. Let me know if you have any cases to add.

In this file, we're just exercising that the deferred diagnostics are working when the kernel lambda is itself templated.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't see any variable template or alias template examples over there. Can you point them out please?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I misunderstood you the first time. I have added alias templates and C++14 variable templates to sycl-restrict.cpp (starting at line 116) as both cases that should be detected and possible false positive cases that should not be flagged.

If you see anything else that we should check, let me know.

class type_info;
typedef __typeof__(sizeof(int)) size_t;
} // namespace std

// testing that the deferred diagnostics work in conjunction with the SYCL namespaces.
inline namespace cl {
namespace sycl {

template <typename name, typename Func>
__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
// expected-note@+1 2{{called by 'kernel_single_task<AName, (lambda}}
// expected-note@+1 3{{called by 'kernel_single_task<AName, (lambda}}
kernelFunc();
}

Expand All @@ -18,6 +23,7 @@ __attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) {
//variadic functions from SYCL kernels emit a deferred diagnostic
void variadic(int, ...) {}

// there are more types like this checked in sycl-restrict.cpp
int calledFromKernel(int a) {
// expected-error@+1 {{zero-length arrays are not permitted in C++}}
int MalArray[0];
Expand All @@ -31,21 +37,102 @@ int calledFromKernel(int a) {
return a + 20;
}

// defines (early and late)
#define floatDef __float128
#define int128Def __int128
#define int128tDef __int128_t
#define intDef int

//typedefs (late )
typedef const __uint128_t megeType;
typedef const __float128 trickyFloatType;
typedef const __int128 tricky128Type;

//templated type (late)
template <typename T>
T bar() { return T(); };

//false positive. early incorrectly catches
template <typename t>
void foo(){};

// template used to specialize a function that contains a lambda that should
// result in a deferred diagnostic being emitted.
// HOWEVER, this is not working presently.
// TODO: re-test after new deferred diagnostic system is merged.
// restore the "FIX!!" tests below

template <typename T>
void setup_sycl_operation(const T VA[]) {

cl::sycl::kernel_single_task<class AName>([]() {
// FIX!! xpected-error@+1 {{zero-length arrays are not permitted in C++}}
int OverlookedBadArray[0];
// ======= Zero Length Arrays Not Allowed in Kernel ==========
// expected-error@+1 {{zero-length arrays are not permitted in C++}}
int MalArray[0];
// expected-error@+1 {{zero-length arrays are not permitted in C++}}
intDef MalArrayDef[0];
// ---- false positive tests. These should not generate any errors.
foo<int[0]>();
std::size_t arrSz = sizeof(int[0]);

// FIX!! xpected-error@+1 {{__float128 is not supported on this target}}
__float128 overlookedBadFloat = 40;
// ======= Float128 Not Allowed in Kernel ==========
// expected-error@+1 {{__float128 is not supported on this target}}
__float128 malFloat = 40;
// expected-error@+1 {{__float128 is not supported on this target}}
trickyFloatType malFloatTrick = 41;
// expected-error@+1 {{__float128 is not supported on this target}}
floatDef malFloatDef = 44;
// expected-error@+1 {{__float128 is not supported on this target}}
auto whatFloat = malFloat;
// expected-error@+1 {{__float128 is not supported on this target}}
auto malAutoTemp5 = bar<__float128>();
// expected-error@+1 {{__float128 is not supported on this target}}
auto malAutoTemp6 = bar<trickyFloatType>();
// expected-error@+1 {{__float128 is not supported on this target}}
decltype(malFloat) malDeclFloat = 42;
// ---- false positive tests
std::size_t someSz = sizeof(__float128);
foo<__float128>();

// ======= __int128 Not Allowed in Kernel ==========
// expected-error@+1 {{__int128 is not supported on this target}}
__int128 malIntent = 2;
// expected-error@+1 {{__int128 is not supported on this target}}
tricky128Type mal128Trick = 2;
// expected-error@+1 {{__int128 is not supported on this target}}
int128Def malIntDef = 9;
// expected-error@+1 {{__int128 is not supported on this target}}
auto whatInt128 = malIntent;
// expected-error@+1 {{__int128 is not supported on this target}}
auto malAutoTemp = bar<__int128>();
// expected-error@+1 {{__int128 is not supported on this target}}
auto malAutoTemp2 = bar<tricky128Type>();
// expected-error@+1 {{__int128 is not supported on this target}}
decltype(malIntent) malDeclInt = 2;

// expected-error@+1 {{__int128 is not supported on this target}}
__int128_t malInt128 = 2;
// expected-error@+1 {{unsigned __int128 is not supported on this target}}
__uint128_t malUInt128 = 3;
// expected-error@+1 {{unsigned __int128 is not supported on this target}}
megeType malTypeDefTrick = 4;
// expected-error@+1 {{__int128 is not supported on this target}}
int128tDef malInt2Def = 6;
// expected-error@+1 {{unsigned __int128 is not supported on this target}}
auto whatUInt = malUInt128;
// expected-error@+1 {{__int128 is not supported on this target}}
auto malAutoTemp3 = bar<__int128_t>();
// expected-error@+1 {{unsigned __int128 is not supported on this target}}
auto malAutoTemp4 = bar<megeType>();
// expected-error@+1 {{__int128 is not supported on this target}}
decltype(malInt128) malDeclInt128 = 5;

// ---- false positive tests These should not generate any errors.
std::size_t i128Sz = sizeof(__int128);
foo<__int128>();
std::size_t u128Sz = sizeof(__uint128_t);
foo<__int128_t>();

// ========= variadic
//expected-error@+1 {{SYCL kernel cannot call a variadic function}}
variadic(5);
});
}

Expand Down
Loading