Skip to content

Commit a09aed0

Browse files
srividya-sundaramsrividya-sundaram
authored
[SYCL] Remove diagnostics emission from Integration header functionality (#2474)
Using Clang's Visitor classes to walk the template type parameter - KernelName to emit the following diagnostics: - Kernel name is invalid. Unscoped enum requires fixed underlying type - enum in template params - Kernel needs to have a globally-visible name - Kernel name is missing
1 parent c3f5cfb commit a09aed0

File tree

5 files changed

+144
-54
lines changed

5 files changed

+144
-54
lines changed

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 108 additions & 40 deletions
Original file line numberDiff line numberDiff line change
@@ -14,6 +14,8 @@
1414
#include "clang/AST/QualTypeNames.h"
1515
#include "clang/AST/RecordLayout.h"
1616
#include "clang/AST/RecursiveASTVisitor.h"
17+
#include "clang/AST/TemplateArgumentVisitor.h"
18+
#include "clang/AST/TypeVisitor.h"
1719
#include "clang/Analysis/CallGraph.h"
1820
#include "clang/Basic/Attributes.h"
1921
#include "clang/Basic/Builtins.h"
@@ -2473,9 +2475,111 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler {
24732475

24742476
} // namespace
24752477

2478+
class SYCLKernelNameTypeVisitor
2479+
: public TypeVisitor<SYCLKernelNameTypeVisitor>,
2480+
public ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor> {
2481+
Sema &S;
2482+
SourceLocation KernelInvocationFuncLoc;
2483+
using InnerTypeVisitor = TypeVisitor<SYCLKernelNameTypeVisitor>;
2484+
using InnerTAVisitor =
2485+
ConstTemplateArgumentVisitor<SYCLKernelNameTypeVisitor>;
2486+
2487+
public:
2488+
SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc)
2489+
: S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc) {}
2490+
2491+
void Visit(QualType T) {
2492+
if (T.isNull())
2493+
return;
2494+
const CXXRecordDecl *RD = T->getAsCXXRecordDecl();
2495+
if (!RD)
2496+
return;
2497+
// If KernelNameType has template args visit each template arg via
2498+
// ConstTemplateArgumentVisitor
2499+
if (const auto *TSD = dyn_cast<ClassTemplateSpecializationDecl>(RD)) {
2500+
const TemplateArgumentList &Args = TSD->getTemplateArgs();
2501+
for (unsigned I = 0; I < Args.size(); I++) {
2502+
Visit(Args[I]);
2503+
}
2504+
} else {
2505+
InnerTypeVisitor::Visit(T.getTypePtr());
2506+
}
2507+
}
2508+
2509+
void Visit(const TemplateArgument &TA) {
2510+
if (TA.isNull())
2511+
return;
2512+
InnerTAVisitor::Visit(TA);
2513+
}
2514+
2515+
void VisitEnumType(const EnumType *T) {
2516+
const EnumDecl *ED = T->getDecl();
2517+
if (!ED->isScoped() && !ED->isFixed()) {
2518+
S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
2519+
<< /* Unscoped enum requires fixed underlying type */ 2;
2520+
S.Diag(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
2521+
<< ED;
2522+
}
2523+
}
2524+
2525+
void VisitRecordType(const RecordType *T) {
2526+
return VisitTagDecl(T->getDecl());
2527+
}
2528+
2529+
void VisitTagDecl(const TagDecl *Tag) {
2530+
bool UnnamedLambdaEnabled =
2531+
S.getASTContext().getLangOpts().SYCLUnnamedLambda;
2532+
if (!Tag->getDeclContext()->isTranslationUnit() &&
2533+
!isa<NamespaceDecl>(Tag->getDeclContext()) && !UnnamedLambdaEnabled) {
2534+
const bool KernelNameIsMissing = Tag->getName().empty();
2535+
if (KernelNameIsMissing) {
2536+
S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named)
2537+
<< /* kernel name is missing */ 0;
2538+
} else {
2539+
if (Tag->isCompleteDefinition())
2540+
S.Diag(KernelInvocationFuncLoc,
2541+
diag::err_sycl_kernel_incorrectly_named)
2542+
<< /* kernel name is not globally-visible */ 1;
2543+
else
2544+
S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl);
2545+
2546+
S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl)
2547+
<< Tag->getName();
2548+
}
2549+
}
2550+
}
2551+
2552+
void VisitTypeTemplateArgument(const TemplateArgument &TA) {
2553+
QualType T = TA.getAsType();
2554+
if (const auto *ET = T->getAs<EnumType>())
2555+
VisitEnumType(ET);
2556+
else
2557+
Visit(T);
2558+
}
2559+
2560+
void VisitIntegralTemplateArgument(const TemplateArgument &TA) {
2561+
QualType T = TA.getIntegralType();
2562+
if (const EnumType *ET = T->getAs<EnumType>())
2563+
VisitEnumType(ET);
2564+
}
2565+
2566+
void VisitTemplateTemplateArgument(const TemplateArgument &TA) {
2567+
TemplateDecl *TD = TA.getAsTemplate().getAsTemplateDecl();
2568+
TemplateParameterList *TemplateParams = TD->getTemplateParameters();
2569+
for (NamedDecl *P : *TemplateParams) {
2570+
if (NonTypeTemplateParmDecl *TemplateParam =
2571+
dyn_cast<NonTypeTemplateParmDecl>(P))
2572+
if (const EnumType *ET = TemplateParam->getType()->getAs<EnumType>())
2573+
VisitEnumType(ET);
2574+
}
2575+
}
2576+
};
2577+
24762578
void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
24772579
ArrayRef<const Expr *> Args) {
24782580
const CXXRecordDecl *KernelObj = getKernelObjectType(KernelFunc);
2581+
QualType KernelNameType =
2582+
calculateKernelNameType(getASTContext(), KernelFunc);
24792583
if (!KernelObj) {
24802584
Diag(Args[0]->getExprLoc(), diag::err_sycl_kernel_not_function_object);
24812585
KernelFunc->setInvalidDecl();
@@ -2511,6 +2615,10 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc,
25112615
return;
25122616

25132617
KernelObjVisitor Visitor{*this};
2618+
SYCLKernelNameTypeVisitor KernelTypeVisitor(*this, Args[0]->getExprLoc());
2619+
// Emit diagnostics for SYCL device kernels only
2620+
if (LangOpts.SYCLIsDevice)
2621+
KernelTypeVisitor.Visit(KernelNameType);
25142622
DiagnosingSYCLKernel = true;
25152623
Visitor.VisitRecordBases(KernelObj, FieldChecker, UnionChecker,
25162624
ArgsSizeChecker);
@@ -2856,18 +2964,6 @@ static void emitWithoutAnonNamespaces(llvm::raw_ostream &OS, StringRef Source) {
28562964
OS << Source;
28572965
}
28582966

2859-
static bool checkEnumTemplateParameter(const EnumDecl *ED,
2860-
DiagnosticsEngine &Diag,
2861-
SourceLocation KernelLocation) {
2862-
if (!ED->isScoped() && !ED->isFixed()) {
2863-
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) << 2;
2864-
Diag.Report(ED->getSourceRange().getBegin(), diag::note_entity_declared_at)
2865-
<< ED;
2866-
return true;
2867-
}
2868-
return false;
2869-
}
2870-
28712967
// Emits a forward declaration
28722968
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
28732969
SourceLocation KernelLocation) {
@@ -2880,32 +2976,6 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
28802976
auto *NS = dyn_cast_or_null<NamespaceDecl>(DC);
28812977

28822978
if (!NS) {
2883-
if (!DC->isTranslationUnit()) {
2884-
const TagDecl *TD = isa<ClassTemplateDecl>(D)
2885-
? cast<ClassTemplateDecl>(D)->getTemplatedDecl()
2886-
: dyn_cast<TagDecl>(D);
2887-
2888-
if (TD && !UnnamedLambdaSupport) {
2889-
// defined class constituting the kernel name is not globally
2890-
// accessible - contradicts the spec
2891-
const bool KernelNameIsMissing = TD->getName().empty();
2892-
if (KernelNameIsMissing) {
2893-
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
2894-
<< /* kernel name is missing */ 0;
2895-
// Don't emit note if kernel name was completely omitted
2896-
} else {
2897-
if (TD->isCompleteDefinition())
2898-
Diag.Report(KernelLocation,
2899-
diag::err_sycl_kernel_incorrectly_named)
2900-
<< /* kernel name is not globally-visible */ 1;
2901-
else
2902-
Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl);
2903-
Diag.Report(D->getSourceRange().getBegin(),
2904-
diag::note_previous_decl)
2905-
<< TD->getName();
2906-
}
2907-
}
2908-
}
29092979
break;
29102980
}
29112981

@@ -3025,7 +3095,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
30253095
// Handle Kernel Name Type templated using enum type and value.
30263096
if (const auto *ET = T->getAs<EnumType>()) {
30273097
const EnumDecl *ED = ET->getDecl();
3028-
if (!checkEnumTemplateParameter(ED, Diag, KernelLocation))
30293098
emitFwdDecl(O, ED, KernelLocation);
30303099
} else if (Arg.getKind() == TemplateArgument::ArgKind::Type)
30313100
emitForwardClassDecls(O, T, KernelLocation, Printed);
@@ -3085,7 +3154,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
30853154
QualType T = TemplateParam->getType();
30863155
if (const auto *ET = T->getAs<EnumType>()) {
30873156
const EnumDecl *ED = ET->getDecl();
3088-
if (!checkEnumTemplateParameter(ED, Diag, KernelLocation))
30893157
emitFwdDecl(O, ED, KernelLocation);
30903158
}
30913159
}

‎clang/test/SemaSYCL/args-size-overflow.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -4,6 +4,7 @@
44
// RUN: %clang_cc1 -fsycl -triple spir64_gen -Werror=sycl-strict -DERROR -fsycl-is-device -fsyntax-only -verify %s
55

66
#include "Inputs/sycl.hpp"
7+
class Foo;
78

89
template <typename Name, typename F>
910
__attribute__((sycl_kernel)) void kernel(F KernelFunc) {
@@ -37,5 +38,5 @@ void use() {
3738
#if defined(GPU) || defined(ERROR)
3839
// expected-note@+2 {{in instantiation of function template specialization 'parallel_for<Foo}}
3940
#endif
40-
parallel_for<class Foo>(L);
41+
parallel_for<Foo>(L);
4142
}

‎clang/test/SemaSYCL/implicit_kernel_type.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -51,10 +51,11 @@ class myWrapper2;
5151
int main() {
5252
queue q;
5353
#ifndef __SYCL_UNNAMED_LAMBDA__
54-
// expected-note@+1 {{InvalidKernelName1 declared here}}
54+
// expected-note@+3 {{InvalidKernelName1 declared here}}
55+
// expected-note@+4{{in instantiation of function template specialization}}
56+
// expected-error@28 {{kernel needs to have a globally-visible name}}
5557
class InvalidKernelName1 {};
5658
q.submit([&](handler &h) {
57-
// expected-error@+1 {{kernel needs to have a globally-visible name}}
5859
h.single_task<InvalidKernelName1>([]() {});
5960
});
6061
#endif

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

Lines changed: 15 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,13 +1,14 @@
11
// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s
22

3+
// expected-error@Inputs/sycl.hpp:220 2{{kernel name is invalid. Unscoped enum requires fixed underlying type}}
34
#include "Inputs/sycl.hpp"
45

56
enum unscoped_enum_int : int {
67
val_1,
78
val_2
89
};
910

10-
// expected-note@+1 {{'unscoped_enum_no_type_set' declared here}}
11+
// expected-note@+1 2 {{'unscoped_enum_no_type_set' declared here}}
1112
enum unscoped_enum_no_type_set {
1213
val_3,
1314
val_4
@@ -29,13 +30,18 @@ class dummy_functor_1 {
2930
void operator()() const {}
3031
};
3132

32-
// expected-error@+2 {{kernel name is invalid. Unscoped enum requires fixed underlying type}}
3333
template <unscoped_enum_no_type_set EnumType>
3434
class dummy_functor_2 {
3535
public:
3636
void operator()() const {}
3737
};
3838

39+
template <template <unscoped_enum_no_type_set EnumType> class C>
40+
class templated_functor {
41+
public:
42+
void operator()() const {}
43+
};
44+
3945
template <scoped_enum_int EnumType>
4046
class dummy_functor_3 {
4147
public:
@@ -54,6 +60,7 @@ int main() {
5460
dummy_functor_2<val_3> f2;
5561
dummy_functor_3<scoped_enum_int::val_2> f3;
5662
dummy_functor_4<scoped_enum_no_type_set::val_4> f4;
63+
templated_functor<dummy_functor_2> f5;
5764

5865
cl::sycl::queue q;
5966

@@ -62,9 +69,15 @@ int main() {
6269
});
6370

6471
q.submit([&](cl::sycl::handler &cgh) {
72+
// expected-note@+1{{in instantiation of function template specialization}}
6573
cgh.single_task(f2);
6674
});
6775

76+
q.submit([&](cl::sycl::handler &cgh) {
77+
// expected-note@+1{{in instantiation of function template specialization}}
78+
cgh.single_task(f5);
79+
});
80+
6881
q.submit([&](cl::sycl::handler &cgh) {
6982
cgh.single_task(f3);
7083
});

‎clang/test/SemaSYCL/unnamed-kernel.cpp

Lines changed: 16 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -27,34 +27,38 @@ struct MyWrapper {
2727
void test() {
2828
cl::sycl::queue q;
2929
#ifndef __SYCL_UNNAMED_LAMBDA__
30-
// expected-error@+5 {{kernel needs to have a globally-visible name}}
31-
// expected-note@+2 {{InvalidKernelName1 declared here}}
30+
// expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
31+
// expected-note@+3 {{InvalidKernelName1 declared here}}
32+
// expected-note@+4{{in instantiation of function template specialization}}
3233
#endif
3334
class InvalidKernelName1 {};
3435
q.submit([&](cl::sycl::handler &h) {
3536
h.single_task<InvalidKernelName1>([] {});
3637
});
3738

3839
#ifndef __SYCL_UNNAMED_LAMBDA__
39-
// expected-error@+5 {{kernel needs to have a globally-visible name}}
40-
// expected-note@+2 {{InvalidKernelName2 declared here}}
40+
// expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
41+
// expected-note@+3 {{InvalidKernelName2 declared here}}
42+
// expected-note@+4{{in instantiation of function template specialization}}
4143
#endif
4244
class InvalidKernelName2 {};
4345
q.submit([&](cl::sycl::handler &h) {
4446
h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
4547
});
4648

4749
#ifndef __SYCL_UNNAMED_LAMBDA__
48-
// expected-error@+4 {{kernel needs to have a globally-visible name}}
50+
// expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
4951
// expected-note@21 {{InvalidKernelName0 declared here}}
52+
// expected-note@+3{{in instantiation of function template specialization}}
5053
#endif
5154
q.submit([&](cl::sycl::handler &h) {
5255
h.single_task<InvalidKernelName0>([] {});
5356
});
5457

5558
#ifndef __SYCL_UNNAMED_LAMBDA__
56-
// expected-error@+4 {{kernel needs to have a globally-visible name}}
59+
// expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
5760
// expected-note@22 {{InvalidKernelName3 declared here}}
61+
// expected-note@+3{{in instantiation of function template specialization}}
5862
#endif
5963
q.submit([&](cl::sycl::handler &h) {
6064
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
@@ -74,17 +78,19 @@ struct MyWrapper {
7478

7579
using InvalidAlias = InvalidKernelName4;
7680
#ifndef __SYCL_UNNAMED_LAMBDA__
77-
// expected-error@+4 {{kernel needs to have a globally-visible name}}
81+
// expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
7882
// expected-note@23 {{InvalidKernelName4 declared here}}
83+
// expected-note@+3{{in instantiation of function template specialization}}
7984
#endif
8085
q.submit([&](cl::sycl::handler &h) {
8186
h.single_task<InvalidAlias>([] {});
8287
});
8388

8489
using InvalidAlias1 = InvalidKernelName5;
8590
#ifndef __SYCL_UNNAMED_LAMBDA__
86-
// expected-error@+4 {{kernel needs to have a globally-visible name}}
91+
// expected-error@Inputs/sycl.hpp:220 {{kernel needs to have a globally-visible name}}
8792
// expected-note@24 {{InvalidKernelName5 declared here}}
93+
// expected-note@+3{{in instantiation of function template specialization}}
8894
#endif
8995
q.submit([&](cl::sycl::handler &h) {
9096
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
@@ -95,7 +101,8 @@ struct MyWrapper {
95101
int main() {
96102
cl::sycl::queue q;
97103
#ifndef __SYCL_UNNAMED_LAMBDA__
98-
// expected-error@+2 {{kernel name is missing}}
104+
// expected-error@Inputs/sycl.hpp:220 {{kernel name is missing}}
105+
// expected-note@+2{{in instantiation of function template specialization}}
99106
#endif
100107
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
101108

0 commit comments

Comments
 (0)