Skip to content

Commit 89886da

Browse files
AlexeySachkovAlexeySachkov
committed
[SYCL] Fix diagnostic about non-globally-visible kernel name
Moved some dead code from CodeGenSYCL/int_header1.cpp to the new test Signed-off-by: Alexey Sachkov <[email protected]>
1 parent 3baec18 commit 89886da

File tree

5 files changed

+120
-76
lines changed

5 files changed

+120
-76
lines changed

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

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -10655,8 +10655,7 @@ def err_builtin_launder_invalid_arg : Error<
1065510655
def err_sycl_attribute_address_space_invalid : Error<
1065610656
"address space is outside the valid range of values">;
1065710657
def err_sycl_kernel_name_class_not_top_level : Error<
10658-
"kernel name class and its template argument classes' declarations can only "
10659-
"nest in a namespace: %0">;
10658+
"kernel needs to have a globally-visible name">;
1066010659
def err_sycl_restrict : Error<
1066110660
"SYCL kernel cannot "
1066210661
"%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: 19 additions & 11 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,12 @@ 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.Report(KernelLocation,
15071508
diag::err_sycl_kernel_name_class_not_top_level);
1509+
if (!TD->getName().empty())
1510+
// Don't emit note if kernel name was completely omitted
1511+
Diag.Report(D->getSourceRange().getBegin(),
1512+
diag::note_previous_decl) << TD->getName();
15081513
}
15091514
}
15101515
break;
@@ -1571,7 +1576,8 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
15711576
// template <typename T1, unsigned int N, typename ...T2> class SimpleVadd;
15721577
//
15731578
void SYCLIntegrationHeader::emitForwardClassDecls(
1574-
raw_ostream &O, QualType T, llvm::SmallPtrSetImpl<const void *> &Printed) {
1579+
raw_ostream &O, QualType T, SourceLocation KernelLocation,
1580+
llvm::SmallPtrSetImpl<const void *> &Printed) {
15751581

15761582
// peel off the pointer types and get the class/struct type:
15771583
for (; T->isPointerType(); T = T->getPointeeType())
@@ -1593,14 +1599,14 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
15931599

15941600
switch (Arg.getKind()) {
15951601
case TemplateArgument::ArgKind::Type:
1596-
emitForwardClassDecls(O, Arg.getAsType(), Printed);
1602+
emitForwardClassDecls(O, Arg.getAsType(), KernelLocation, Printed);
15971603
break;
15981604
case TemplateArgument::ArgKind::Pack: {
15991605
ArrayRef<TemplateArgument> Pack = Arg.getPackAsArray();
16001606

16011607
for (const auto &T : Pack) {
16021608
if (T.getKind() == TemplateArgument::ArgKind::Type) {
1603-
emitForwardClassDecls(O, T.getAsType(), Printed);
1609+
emitForwardClassDecls(O, T.getAsType(), KernelLocation, Printed);
16041610
}
16051611
}
16061612
break;
@@ -1623,7 +1629,7 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
16231629
// class template <template <typename> class> class T as it should.
16241630
TemplateDecl *TD = Arg.getAsTemplate().getAsTemplateDecl();
16251631
if (Printed.insert(TD).second) {
1626-
emitFwdDecl(O, TD);
1632+
emitFwdDecl(O, TD, KernelLocation);
16271633
}
16281634
break;
16291635
}
@@ -1637,12 +1643,12 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
16371643
assert(CTD && "template declaration must be available");
16381644

16391645
if (Printed.insert(CTD).second) {
1640-
emitFwdDecl(O, CTD);
1646+
emitFwdDecl(O, CTD, KernelLocation);
16411647
}
16421648
} else if (Printed.insert(RD).second) {
16431649
// emit forward declarations for "leaf" classes in the template parameter
16441650
// tree;
1645-
emitFwdDecl(O, RD);
1651+
emitFwdDecl(O, RD, KernelLocation);
16461652
}
16471653
}
16481654

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

16601666
llvm::SmallPtrSet<const void *, 4> Printed;
16611667
for (const KernelDesc &K : KernelDescs) {
1662-
emitForwardClassDecls(O, K.NameType, Printed);
1668+
emitForwardClassDecls(O, K.NameType, K.KernelLocation, Printed);
16631669
}
16641670
}
16651671
O << "\n";
@@ -1779,11 +1785,13 @@ bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) {
17791785

17801786
void SYCLIntegrationHeader::startKernel(StringRef KernelName,
17811787
QualType KernelNameType,
1782-
StringRef KernelStableName) {
1788+
StringRef KernelStableName,
1789+
SourceLocation KernelLocation) {
17831790
KernelDescs.resize(KernelDescs.size() + 1);
17841791
KernelDescs.back().Name = std::string(KernelName);
17851792
KernelDescs.back().NameType = KernelNameType;
17861793
KernelDescs.back().StableName = std::string(KernelStableName);
1794+
KernelDescs.back().KernelLocation = KernelLocation;
17871795
}
17881796

17891797
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: 90 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,90 @@
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> class KernelName;
11+
}
12+
13+
struct MyWrapper {
14+
private:
15+
class InvalidKernelName0 {};
16+
class InvalidKernelName3 {};
17+
class InvalidKernelName4 {};
18+
class InvalidKernelName5 {};
19+
20+
public:
21+
void test() {
22+
cl::sycl::queue q;
23+
#ifndef __SYCL_UNNAMED_LAMBDA__
24+
// expected-error@+5 {{kernel needs to have a globally-visible name}}
25+
// expected-note@+2 {{InvalidKernelName1 declared here}}
26+
#endif
27+
class InvalidKernelName1 {};
28+
q.submit([&](cl::sycl::handler &h) {
29+
h.single_task<InvalidKernelName1>([] {});
30+
});
31+
32+
#ifndef __SYCL_UNNAMED_LAMBDA__
33+
// expected-error@+5 {{kernel needs to have a globally-visible name}}
34+
// expected-note@+2 {{InvalidKernelName2 declared here}}
35+
#endif
36+
class InvalidKernelName2 {};
37+
q.submit([&](cl::sycl::handler &h) {
38+
h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
39+
});
40+
41+
#ifndef __SYCL_UNNAMED_LAMBDA__
42+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
43+
// expected-note@15 {{InvalidKernelName0 declared here}}
44+
#endif
45+
q.submit([&](cl::sycl::handler &h) {
46+
h.single_task<InvalidKernelName0>([] {});
47+
});
48+
49+
#ifndef __SYCL_UNNAMED_LAMBDA__
50+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
51+
// expected-note@16 {{InvalidKernelName3 declared here}}
52+
#endif
53+
q.submit([&](cl::sycl::handler &h) {
54+
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
55+
});
56+
57+
using ValidAlias = MyWrapper;
58+
q.submit([&](cl::sycl::handler &h) {
59+
h.single_task<ValidAlias>([] {});
60+
});
61+
62+
using InvalidAlias = InvalidKernelName4;
63+
#ifndef __SYCL_UNNAMED_LAMBDA__
64+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
65+
// expected-note@17 {{InvalidKernelName4 declared here}}
66+
#endif
67+
q.submit([&](cl::sycl::handler &h) {
68+
h.single_task<InvalidAlias>([] {});
69+
});
70+
71+
using InvalidAlias1 = InvalidKernelName5;
72+
#ifndef __SYCL_UNNAMED_LAMBDA__
73+
// expected-error@+4 {{kernel needs to have a globally-visible name}}
74+
// expected-note@18 {{InvalidKernelName5 declared here}}
75+
#endif
76+
q.submit([&](cl::sycl::handler &h) {
77+
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
78+
});
79+
}
80+
};
81+
82+
int main() {
83+
cl::sycl::queue q;
84+
#ifndef __SYCL_UNNAMED_LAMBDA__
85+
// expected-error@+2 {{kernel needs to have a globally-visible name}}
86+
#endif
87+
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });
88+
89+
return 0;
90+
}

0 commit comments

Comments
 (0)