Skip to content

Commit fadaa59

Browse files
AlexeySachkovAlexeySachkov
authored
[SYCL] Fix diagnostic about non-globally-visible kernel name (#928)
Moved some dead code from CodeGenSYCL/int_header1.cpp to the new test Improve diagnostic about invalid kernel name Signed-off-by: Alexey Sachkov <[email protected]>
1 parent 4a2aa28 commit fadaa59

File tree

5 files changed

+130
-78
lines changed

5 files changed

+130
-78
lines changed

‎clang/include/clang/Basic/DiagnosticSemaKinds.td

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -10650,9 +10650,9 @@ def err_builtin_launder_invalid_arg : Error<
1065010650
// SYCL-specific diagnostics
1065110651
def err_sycl_attribute_address_space_invalid : Error<
1065210652
"address space is outside the valid range of values">;
10653-
def err_sycl_kernel_name_class_not_top_level : Error<
10654-
"kernel name class and its template argument classes' declarations can only "
10655-
"nest in a namespace: %0">;
10653+
def err_sycl_kernel_incorrectly_named : Error<
10654+
"kernel %select{name is missing"
10655+
"|needs to have a globally-visible name}0">;
1065610656
def err_sycl_restrict : Error<
1065710657
"SYCL kernel cannot "
1065810658
"%select{use a non-const global variable"

‎clang/include/clang/Sema/Sema.h

Lines changed: 10 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -329,7 +329,7 @@ class SYCLIntegrationHeader {
329329
/// Signals that subsequent parameter descriptor additions will go to
330330
/// the kernel with given name. Starts new kernel invocation descriptor.
331331
void startKernel(StringRef KernelName, QualType KernelNameType,
332-
StringRef KernelStableName);
332+
StringRef KernelStableName, SourceLocation Loc);
333333

334334
/// Adds a kernel parameter descriptor to current kernel invocation
335335
/// descriptor.
@@ -367,6 +367,8 @@ class SYCLIntegrationHeader {
367367
/// Kernel name with stable lambda name mangling
368368
std::string StableName;
369369

370+
SourceLocation KernelLocation;
371+
370372
/// Descriptor of kernel actual parameters.
371373
SmallVector<KernelParamDesc, 8> Params;
372374

@@ -381,7 +383,8 @@ class SYCLIntegrationHeader {
381383
}
382384

383385
/// Emits a forward declaration for given declaration.
384-
void emitFwdDecl(raw_ostream &O, const Decl *D);
386+
void emitFwdDecl(raw_ostream &O, const Decl *D,
387+
SourceLocation KernelLocation);
385388

386389
/// Emits forward declarations of classes and template classes on which
387390
/// declaration of given type depends. See example in the comments for the
@@ -390,10 +393,14 @@ class SYCLIntegrationHeader {
390393
/// stream to emit to
391394
/// \param T
392395
/// type to emit forward declarations for
396+
/// \param KernelLocation
397+
/// source location of the SYCL kernel function, used to emit nicer
398+
/// diagnostic messages if kernel name is missing
393399
/// \param Emitted
394400
/// a set of declarations forward declrations has been emitted for already
395401
void emitForwardClassDecls(raw_ostream &O, QualType T,
396-
llvm::SmallPtrSetImpl<const void*> &Emitted);
402+
SourceLocation KernelLocation,
403+
llvm::SmallPtrSetImpl<const void *> &Emitted);
397404

398405
private:
399406
/// Keeps invocation descriptors for each kernel invocation started by

‎clang/lib/Sema/SemaSYCL.cpp

Lines changed: 26 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -1056,7 +1056,7 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
10561056
const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(KernelObjTy);
10571057
const std::string StableName = PredefinedExpr::ComputeName(
10581058
Ctx, PredefinedExpr::UniqueStableNameExpr, NameType);
1059-
H.startKernel(Name, NameType, StableName);
1059+
H.startKernel(Name, NameType, StableName, KernelObjTy->getLocation());
10601060

10611061
auto populateHeaderForAccessor = [&](const QualType &ArgTy, uint64_t Offset) {
10621062
// The parameter is a SYCL accessor object.
@@ -1485,7 +1485,8 @@ static std::string eraseAnonNamespace(std::string S) {
14851485
}
14861486

14871487
// Emits a forward declaration
1488-
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
1488+
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
1489+
SourceLocation KernelLocation) {
14891490
// wrap the declaration into namespaces if needed
14901491
unsigned NamespaceCnt = 0;
14911492
std::string NSStr = "";
@@ -1503,8 +1504,18 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
15031504
if (TD && TD->isCompleteDefinition() && !UnnamedLambdaSupport) {
15041505
// defined class constituting the kernel name is not globally
15051506
// accessible - contradicts the spec
1506-
Diag.Report(D->getSourceRange().getBegin(),
1507-
diag::err_sycl_kernel_name_class_not_top_level);
1507+
const bool KernelNameIsMissing = TD->getName().empty();
1508+
if (KernelNameIsMissing) {
1509+
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
1510+
<< /* kernel name is missing */ 0;
1511+
// Don't emit note if kernel name was completely omitted
1512+
} else {
1513+
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
1514+
<< /* kernel name is not globally-visible */ 1;
1515+
Diag.Report(D->getSourceRange().getBegin(),
1516+
diag::note_previous_decl)
1517+
<< TD->getName();
1518+
}
15081519
}
15091520
}
15101521
break;
@@ -1571,7 +1582,8 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
15711582
// template <typename T1, unsigned int N, typename ...T2> class SimpleVadd;
15721583
//
15731584
void SYCLIntegrationHeader::emitForwardClassDecls(
1574-
raw_ostream &O, QualType T, llvm::SmallPtrSetImpl<const void *> &Printed) {
1585+
raw_ostream &O, QualType T, SourceLocation KernelLocation,
1586+
llvm::SmallPtrSetImpl<const void *> &Printed) {
15751587

15761588
// peel off the pointer types and get the class/struct type:
15771589
for (; T->isPointerType(); T = T->getPointeeType())
@@ -1593,14 +1605,14 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
15931605

15941606
switch (Arg.getKind()) {
15951607
case TemplateArgument::ArgKind::Type:
1596-
emitForwardClassDecls(O, Arg.getAsType(), Printed);
1608+
emitForwardClassDecls(O, Arg.getAsType(), KernelLocation, Printed);
15971609
break;
15981610
case TemplateArgument::ArgKind::Pack: {
15991611
ArrayRef<TemplateArgument> Pack = Arg.getPackAsArray();
16001612

16011613
for (const auto &T : Pack) {
16021614
if (T.getKind() == TemplateArgument::ArgKind::Type) {
1603-
emitForwardClassDecls(O, T.getAsType(), Printed);
1615+
emitForwardClassDecls(O, T.getAsType(), KernelLocation, Printed);
16041616
}
16051617
}
16061618
break;
@@ -1623,7 +1635,7 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
16231635
// class template <template <typename> class> class T as it should.
16241636
TemplateDecl *TD = Arg.getAsTemplate().getAsTemplateDecl();
16251637
if (Printed.insert(TD).second) {
1626-
emitFwdDecl(O, TD);
1638+
emitFwdDecl(O, TD, KernelLocation);
16271639
}
16281640
break;
16291641
}
@@ -1637,12 +1649,12 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
16371649
assert(CTD && "template declaration must be available");
16381650

16391651
if (Printed.insert(CTD).second) {
1640-
emitFwdDecl(O, CTD);
1652+
emitFwdDecl(O, CTD, KernelLocation);
16411653
}
16421654
} else if (Printed.insert(RD).second) {
16431655
// emit forward declarations for "leaf" classes in the template parameter
16441656
// tree;
1645-
emitFwdDecl(O, RD);
1657+
emitFwdDecl(O, RD, KernelLocation);
16461658
}
16471659
}
16481660

@@ -1659,7 +1671,7 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) {
16591671

16601672
llvm::SmallPtrSet<const void *, 4> Printed;
16611673
for (const KernelDesc &K : KernelDescs) {
1662-
emitForwardClassDecls(O, K.NameType, Printed);
1674+
emitForwardClassDecls(O, K.NameType, K.KernelLocation, Printed);
16631675
}
16641676
}
16651677
O << "\n";
@@ -1779,11 +1791,13 @@ bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) {
17791791

17801792
void SYCLIntegrationHeader::startKernel(StringRef KernelName,
17811793
QualType KernelNameType,
1782-
StringRef KernelStableName) {
1794+
StringRef KernelStableName,
1795+
SourceLocation KernelLocation) {
17831796
KernelDescs.resize(KernelDescs.size() + 1);
17841797
KernelDescs.back().Name = std::string(KernelName);
17851798
KernelDescs.back().NameType = KernelNameType;
17861799
KernelDescs.back().StableName = std::string(KernelStableName);
1800+
KernelDescs.back().KernelLocation = KernelLocation;
17871801
}
17881802

17891803
void SYCLIntegrationHeader::addParamDesc(kernel_param_kind_t Kind, int Info,

‎clang/test/CodeGenSYCL/int_header1.cpp

Lines changed: 0 additions & 60 deletions
Original file line numberDiff line numberDiff line change
@@ -76,22 +76,6 @@ struct MyWrapper {
7676
// translation unit scope
7777
kernel_single_task<nm1::nm2::KernelName0>([=]() { acc.use(); });
7878

79-
#ifdef LI__
80-
// TODO unexpected compilation error when host code + integration header
81-
// is compiled LI-- kernel name is an incomplete class forward-declared in
82-
// local scope
83-
class KernelName2;
84-
kernel_single_task<KernelName2>([=]() { acc.use(); });
85-
#endif
86-
87-
#ifdef LD__
88-
// Expected compilation error.
89-
// LD--
90-
// kernel name is a class defined in local scope
91-
class KernelName2a {};
92-
kernel_single_task<KernelName2a>([=]() { acc.use(); });
93-
#endif
94-
9579
// TI--
9680
// an incomplete class forward-declared in a namespace at
9781
// translation unit scope
@@ -127,16 +111,6 @@ struct MyWrapper {
127111
kernel_single_task<nm1::KernelName3<class KernelName5>>(
128112
[=]() { acc.use(); });
129113

130-
#ifdef TILI
131-
// Expected compilation error
132-
// TILI
133-
// an incomplete template specialization class with incomplete class as
134-
// argument forward-declared locally
135-
class KernelName6;
136-
kernel_single_task<nm1::KernelName3<KernelName6>>(
137-
[=]() { acc.use(); });
138-
#endif
139-
140114
// TDPI
141115
// a defined template specialization class with incomplete class as
142116
// argument forward-declared "in-place"
@@ -149,40 +123,6 @@ struct MyWrapper {
149123
kernel_single_task<nm1::KernelName8<nm1::nm2::C>>(
150124
[=]() { acc.use(); });
151125

152-
#ifdef TDLI
153-
// TODO unexpected compilation error when host code + integration header
154-
// is compiled TDLI a defined template specialization class with
155-
// incomplete class as argument forward-declared locally
156-
class KernelName6a;
157-
kernel_single_task<nm1::KernelName4<KernelName6a>>(
158-
[=]() { acc.use(); });
159-
#endif
160-
161-
#ifdef TDLD
162-
// Expected compilation error
163-
// TDLD
164-
// a defined template specialization class with a class as argument
165-
// defined locally
166-
class KernelName9 {};
167-
kernel_single_task<nm1::KernelName4<KernelName9>>([=]() { acc.use(); });
168-
#endif
169-
170-
#ifdef TICD
171-
// Expected compilation error
172-
// TICD
173-
// an incomplete template specialization class with a defined class as
174-
// argument declared in the containing class
175-
kernel_single_task<nm1::KernelName3<KN100>>([=]() { acc.use(); });
176-
#endif
177-
178-
#ifdef TICI
179-
// Expected compilation error
180-
// TICI
181-
// an incomplete template specialization class with an incomplete class as
182-
// argument declared in the containing class
183-
kernel_single_task<nm1::KernelName3<KN101>>([=]() { acc.use(); });
184-
#endif
185-
186126
// kernel name type is a templated class, both the top-level class and the
187127
// template argument are declared in the anonymous namespace
188128
kernel_single_task<TmplClassInAnonNS<class ClassInAnonNS>>(

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

Lines changed: 91 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,91 @@
1+
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
2+
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s
3+
#include <sycl.hpp>
4+
5+
#ifdef __SYCL_UNNAMED_LAMBDA__
6+
// expected-no-diagnostics
7+
#endif
8+
9+
namespace namespace1 {
10+
template <typename T>
11+
class KernelName;
12+
}
13+
14+
struct MyWrapper {
15+
private:
16+
class InvalidKernelName0 {};
17+
class InvalidKernelName3 {};
18+
class InvalidKernelName4 {};
19+
class InvalidKernelName5 {};
20+
21+
public:
22+
void test() {
23+
cl::sycl::queue q;
24+
#ifndef __SYCL_UNNAMED_LAMBDA__
25+
// expected-error@+5 {{kernel needs to have a globally-visible name}}
26+
// expected-note@+2 {{InvalidKernelName1 declared here}}
27+
#endif
28+
class InvalidKernelName1 {};
29+
q.submit([&](cl::sycl::handler &h) {
30+
h.single_task<InvalidKernelName1>([] {});
31+
});
32+
33+
#ifndef __SYCL_UNNAMED_LAMBDA__
34+
// expected-error@+5 {{kernel needs to have a globally-visible name}}
35+
// expected-note@+2 {{InvalidKernelName2 declared here}}
36+
#endif
37+
class InvalidKernelName2 {};
38+
q.submit([&](cl::sycl::handler &h) {
39+
h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
40+
});
41+
42+
#ifndef __SYCL_UNNAMED_LAMBDA__
43+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
44+
// expected-note@16 {{InvalidKernelName0 declared here}}
45+
#endif
46+
q.submit([&](cl::sycl::handler &h) {
47+
h.single_task<InvalidKernelName0>([] {});
48+
});
49+
50+
#ifndef __SYCL_UNNAMED_LAMBDA__
51+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
52+
// expected-note@17 {{InvalidKernelName3 declared here}}
53+
#endif
54+
q.submit([&](cl::sycl::handler &h) {
55+
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
56+
});
57+
58+
using ValidAlias = MyWrapper;
59+
q.submit([&](cl::sycl::handler &h) {
60+
h.single_task<ValidAlias>([] {});
61+
});
62+
63+
using InvalidAlias = InvalidKernelName4;
64+
#ifndef __SYCL_UNNAMED_LAMBDA__
65+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
66+
// expected-note@18 {{InvalidKernelName4 declared here}}
67+
#endif
68+
q.submit([&](cl::sycl::handler &h) {
69+
h.single_task<InvalidAlias>([] {});
70+
});
71+
72+
using InvalidAlias1 = InvalidKernelName5;
73+
#ifndef __SYCL_UNNAMED_LAMBDA__
74+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
75+
// expected-note@19 {{InvalidKernelName5 declared here}}
76+
#endif
77+
q.submit([&](cl::sycl::handler &h) {
78+
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
79+
});
80+
}
81+
};
82+
83+
int main() {
84+
cl::sycl::queue q;
85+
#ifndef __SYCL_UNNAMED_LAMBDA__
86+
// expected-error@+2 {{kernel name is missing}}
87+
#endif
88+
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
89+
90+
return 0;
91+
}

0 commit comments

Comments
 (0)