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 13 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 checkSYCLDeviceVarDecl(VarDecl *Var);
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)
checkSYCLDeviceVarDecl(var);

// In Objective-C, don't allow jumps past the implicit initialization of a
// local retaining variable.
if (getLangOpts().ObjC &&
Expand Down
90 changes: 90 additions & 0 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,96 @@ bool Sema::isKnownGoodSYCLDecl(const Decl *D) {
return false;
}

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

static Sema::DeviceDiagBuilder
emitDeferredDiagnosticAndNote(Sema &S, SourceRange Loc, unsigned DiagID,
SourceRange UsedAtLoc) {
Sema::DeviceDiagBuilder builder =
S.SYCLDiagIfDeviceCode(Loc.getBegin(), DiagID);
if (UsedAtLoc.isValid())
S.SYCLDiagIfDeviceCode(UsedAtLoc.getBegin(), diag::note_sycl_used_here);
return builder;
}

static void checkSYCLVarType(Sema &S, QualType Ty, SourceRange Loc,
llvm::DenseSet<QualType> Visited,
SourceRange UsedAtLoc = SourceRange()) {
// Not all variable types are supported inside SYCL kernels,
// for example the quad type __float128 will cause the resulting
// SPIR-V to not link.
Copy link
Contributor

Choose a reason for hiding this comment

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

__float128 will cause the resulting SPIR-V to not link.

What exactly do you mean here? AFAIK we don't link SPIR-V files.

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 am not an expert on the pipeline, especially in regards to how SPIR-V is handled. In the simplest case, without these type checks, if you have kernel code that contains an unsupported type (such as __int128) , you can pass the -c flag to the compiler to simply output the object files, and everything is fine. But when you next go to link those files into your executable, you'll get an error ( InvalidBitWidth: Invalid bit width in input: 128 etc) . IIRC, those are emitted by SpirVWriter.cpp
So it may not be linking, per se. But it occurs in what we normally think of as the link phase.

Copy link
Contributor

Choose a reason for hiding this comment

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

Work with @Fznamznon to make sure that this comment means what you want it to.

Copy link
Contributor

@Fznamznon Fznamznon Apr 7, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

Hmm okay, such errors actually happen during LLVM IR -> SPIR-V translation and it really happens somewhere near the link phase, but technically it is not linking.
Let's do not confuse anyone and mention something like that: "unsupported types in device code cause errors on SPIR-V translation stage." You can rewrite it in a better way, but please keep it true.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

updated it

// Here we check any potentially unsupported declaration and issue
// a deferred diagnostic, which will be emitted iff the declaration
// is discovered to reside in kernel code.
// The optional UsedAtLoc param is used when the SYCL usage is at a
// different location than the variable declaration and we need to
// inform the user of both, e.g. struct member usage vs declaration.

//--- check types ---

// zero length arrays
if (isZeroSizedArray(Ty))
emitDeferredDiagnosticAndNote(S, Loc, diag::err_typecheck_zero_array_size,
UsedAtLoc);

// Sub-reference array or pointer, then proceed with that type.
while (Ty->isAnyPointerType() || Ty->isArrayType())
Ty = QualType{Ty->getPointeeOrArrayElementType(), 0};

// __int128, __int128_t, __uint128_t, __float128
if (Ty->isSpecificBuiltinType(BuiltinType::Int128) ||
Ty->isSpecificBuiltinType(BuiltinType::UInt128) ||
(Ty->isSpecificBuiltinType(BuiltinType::Float128) &&
!S.Context.getTargetInfo().hasFloat128Type()))
emitDeferredDiagnosticAndNote(S, Loc, diag::err_type_unsupported, UsedAtLoc)
<< Ty.getUnqualifiedType().getCanonicalType();

//--- now recurse ---
// Pointers complicate recursion. Add this type to Visited.
// If already there, bail out.
if (!Visited.insert(Ty).second)
return;

if (const auto *ATy = dyn_cast<AttributedType>(Ty))
return checkSYCLVarType(S, ATy->getModifiedType(), Loc, Visited);

if (const auto *CRD = Ty->getAsCXXRecordDecl()) {
// If the class is a forward declaration - skip it, because otherwise we
// would query property of class with no definition, which results in
// clang crash.
Copy link
Contributor

Choose a reason for hiding this comment

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

Which thing are you querying that would cause this? From the rest of this condition, you simply go through 'fields', which would just be empty if there wasn't a definition, right?

Copy link
Contributor Author

Choose a reason for hiding this comment

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

This recursion code, and its comments, was brought over from the recursion in the CheckSYCLType routine that is called by the AST Visitors. I was hoping to leverage the experience there.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

But I'll see if I this is strictly necessary here and remove it if not.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

removed

if (!CRD->hasDefinition())
return;

for (const auto &Field : CRD->fields())
checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited,
Loc);
} else if (const auto *RD = Ty->getAsRecordDecl()) {
for (const auto &Field : RD->fields())
checkSYCLVarType(S, Field->getType(), Field->getSourceRange(), Visited,
Loc);
} else if (const auto *FPTy = dyn_cast<FunctionProtoType>(Ty)) {
for (const auto &ParamTy : FPTy->param_types())
checkSYCLVarType(S, ParamTy, Loc, Visited);
checkSYCLVarType(S, FPTy->getReturnType(), Loc, Visited);
} else if (const auto *FTy = dyn_cast<FunctionType>(Ty)) {
Copy link
Contributor

Choose a reason for hiding this comment

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

Again, curious as to where you found a function type without a prototype in C++ mode! You cannot actually spell one in SYCL/C++. I'm not sure this else-if is worth having.

Copy link
Contributor Author

@cperkinsintel cperkinsintel Apr 7, 2020 •
edited
Loading

Choose a reason for hiding this comment

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

@erichkeane in retrospect, I'm not sure that any of these additional paths are needed.

If this recursive part was reduced to simply running through the fields of a CXXRecordDecl, and nothing more, would we be overlooking anything? I've been testing function declarations, function pointers, extern "C" structs, etc. and none of them require these paths.

Copy link
Contributor

Choose a reason for hiding this comment

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

FYI: Just loop through the fields of a RecordDecl (not CXXRecordDecl). Its a little 'shorter' and makes the transition to allowing this code someday in C code lighter (as well as being the 'lowest' type that has the fields you need).

I'm quite surprised the function type doesn't go through there. The CheckSYCLType definitely gets FunctionProtoType going through there(just auto V = &func;), so I'd expect us to have to check that.

FunctionType should NOT get hit, since it isn't possible to create a FunctionType that isn't also a FunctionProtoType[0] (or RecordDecl (that isn't also a CXXRecordDecl) for that matter) in C++.

Basically, Only in C mode do we create RecordDecl (again, that isn't a CXXRecordDecl). Also, it is only possible to spell a Function without a prototype in C. For example:

void foo();

In C, this is a function without a prototype, we don't know the parameter list. In C++ DOES have a prototype, and is equivalent to:

void foo(void);

Copy link
Contributor Author

Choose a reason for hiding this comment

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

thanks. that helped, and gave me a new test case as well.
I've tightened up the recursion code considerably, and added a test that verifies function prototypes are getting these type checks applied.

checkSYCLVarType(S, FTy->getReturnType(), Loc, Visited);
}
}

void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) {
assert(getLangOpts().SYCLIsDevice &&
"Should only be called during SYCL compilation");
QualType Ty = Var->getType();
SourceRange Loc = Var->getLocation();
llvm::DenseSet<QualType> Visited;

checkSYCLVarType(*this, Ty, Loc, Visited);
}

class MarkDeviceFunction : public RecursiveASTVisitor<MarkDeviceFunction> {
public:
MarkDeviceFunction(Sema &S)
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
109 changes: 98 additions & 11 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,11 +23,12 @@ __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];

// expected-error@+1 {{__float128 is not supported on this target}}
// expected-error@+1 {{'__float128' is not supported on this target}}
__float128 malFloat = 40;

//expected-error@+1 {{SYCL kernel cannot call a variadic function}}
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];

// FIX!! xpected-error@+1 {{__float128 is not supported on this target}}
__float128 overlookedBadFloat = 40;
// ======= 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]);

// ======= 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 All @@ -56,7 +143,7 @@ int main(int argc, char **argv) {
// expected-error@+1 {{zero-length arrays are not permitted in C++}}
int BadArray[0];

// expected-error@+1 {{__float128 is not supported on this target}}
// expected-error@+1 {{'__float128' is not supported on this target}}
__float128 badFloat = 40; // this SHOULD trigger a diagnostic

//expected-error@+1 {{SYCL kernel cannot call a variadic function}}
Expand Down
Loading