Skip to content

[SYCL] Fix diagnostic about non-globally-visible kernel name #928

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.

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
Show file tree
Hide file tree
Changes from all commits
Commits
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
6 changes: 3 additions & 3 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
Original file line number Diff line number Diff line change
Expand Up @@ -10654,9 +10654,9 @@ def err_builtin_launder_invalid_arg : Error<
// SYCL-specific diagnostics
def err_sycl_attribute_address_space_invalid : Error<
"address space is outside the valid range of values">;
def err_sycl_kernel_name_class_not_top_level : Error<
"kernel name class and its template argument classes' declarations can only "
"nest in a namespace: %0">;
def err_sycl_kernel_incorrectly_named : Error<
"kernel %select{name is missing"
"|needs to have a globally-visible name}0">;
def err_sycl_restrict : Error<
"SYCL kernel cannot "
"%select{use a non-const global variable"
Expand Down
13 changes: 10 additions & 3 deletions clang/include/clang/Sema/Sema.h
Original file line number Diff line number Diff line change
Expand Up @@ -329,7 +329,7 @@ class SYCLIntegrationHeader {
/// Signals that subsequent parameter descriptor additions will go to
/// the kernel with given name. Starts new kernel invocation descriptor.
void startKernel(StringRef KernelName, QualType KernelNameType,
StringRef KernelStableName);
StringRef KernelStableName, SourceLocation Loc);

/// Adds a kernel parameter descriptor to current kernel invocation
/// descriptor.
Expand Down Expand Up @@ -367,6 +367,8 @@ class SYCLIntegrationHeader {
/// Kernel name with stable lambda name mangling
std::string StableName;

SourceLocation KernelLocation;

/// Descriptor of kernel actual parameters.
SmallVector<KernelParamDesc, 8> Params;

Expand All @@ -381,7 +383,8 @@ class SYCLIntegrationHeader {
}

/// Emits a forward declaration for given declaration.
void emitFwdDecl(raw_ostream &O, const Decl *D);
void emitFwdDecl(raw_ostream &O, const Decl *D,
SourceLocation KernelLocation);

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

private:
/// Keeps invocation descriptors for each kernel invocation started by
Expand Down
38 changes: 26 additions & 12 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1056,7 +1056,7 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name,
const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(KernelObjTy);
const std::string StableName = PredefinedExpr::ComputeName(
Ctx, PredefinedExpr::UniqueStableNameExpr, NameType);
H.startKernel(Name, NameType, StableName);
H.startKernel(Name, NameType, StableName, KernelObjTy->getLocation());

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

// Emits a forward declaration
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D,
SourceLocation KernelLocation) {
// wrap the declaration into namespaces if needed
unsigned NamespaceCnt = 0;
std::string NSStr = "";
Expand All @@ -1503,8 +1504,18 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
if (TD && TD->isCompleteDefinition() && !UnnamedLambdaSupport) {
// defined class constituting the kernel name is not globally
// accessible - contradicts the spec
Diag.Report(D->getSourceRange().getBegin(),
diag::err_sycl_kernel_name_class_not_top_level);
const bool KernelNameIsMissing = TD->getName().empty();
if (KernelNameIsMissing) {
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name is missing */ 0;
// Don't emit note if kernel name was completely omitted
} else {
Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named)
<< /* kernel name is not globally-visible */ 1;
Diag.Report(D->getSourceRange().getBegin(),
diag::note_previous_decl)
<< TD->getName();
}
}
}
break;
Expand Down Expand Up @@ -1571,7 +1582,8 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) {
// template <typename T1, unsigned int N, typename ...T2> class SimpleVadd;
//
void SYCLIntegrationHeader::emitForwardClassDecls(
raw_ostream &O, QualType T, llvm::SmallPtrSetImpl<const void *> &Printed) {
raw_ostream &O, QualType T, SourceLocation KernelLocation,
llvm::SmallPtrSetImpl<const void *> &Printed) {

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

switch (Arg.getKind()) {
case TemplateArgument::ArgKind::Type:
emitForwardClassDecls(O, Arg.getAsType(), Printed);
emitForwardClassDecls(O, Arg.getAsType(), KernelLocation, Printed);
break;
case TemplateArgument::ArgKind::Pack: {
ArrayRef<TemplateArgument> Pack = Arg.getPackAsArray();

for (const auto &T : Pack) {
if (T.getKind() == TemplateArgument::ArgKind::Type) {
emitForwardClassDecls(O, T.getAsType(), Printed);
emitForwardClassDecls(O, T.getAsType(), KernelLocation, Printed);
}
}
break;
Expand All @@ -1623,7 +1635,7 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
// class template <template <typename> class> class T as it should.
TemplateDecl *TD = Arg.getAsTemplate().getAsTemplateDecl();
if (Printed.insert(TD).second) {
emitFwdDecl(O, TD);
emitFwdDecl(O, TD, KernelLocation);
}
break;
}
Expand All @@ -1637,12 +1649,12 @@ void SYCLIntegrationHeader::emitForwardClassDecls(
assert(CTD && "template declaration must be available");

if (Printed.insert(CTD).second) {
emitFwdDecl(O, CTD);
emitFwdDecl(O, CTD, KernelLocation);
}
} else if (Printed.insert(RD).second) {
// emit forward declarations for "leaf" classes in the template parameter
// tree;
emitFwdDecl(O, RD);
emitFwdDecl(O, RD, KernelLocation);
}
}

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

llvm::SmallPtrSet<const void *, 4> Printed;
for (const KernelDesc &K : KernelDescs) {
emitForwardClassDecls(O, K.NameType, Printed);
emitForwardClassDecls(O, K.NameType, K.KernelLocation, Printed);
}
}
O << "\n";
Expand Down Expand Up @@ -1779,11 +1791,13 @@ bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) {

void SYCLIntegrationHeader::startKernel(StringRef KernelName,
QualType KernelNameType,
StringRef KernelStableName) {
StringRef KernelStableName,
SourceLocation KernelLocation) {
KernelDescs.resize(KernelDescs.size() + 1);
KernelDescs.back().Name = std::string(KernelName);
KernelDescs.back().NameType = KernelNameType;
KernelDescs.back().StableName = std::string(KernelStableName);
KernelDescs.back().KernelLocation = KernelLocation;
}

void SYCLIntegrationHeader::addParamDesc(kernel_param_kind_t Kind, int Info,
Expand Down
60 changes: 0 additions & 60 deletions clang/test/CodeGenSYCL/int_header1.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -76,22 +76,6 @@ struct MyWrapper {
// translation unit scope
kernel_single_task<nm1::nm2::KernelName0>([=]() { acc.use(); });

#ifdef LI__
// TODO unexpected compilation error when host code + integration header
// is compiled LI-- kernel name is an incomplete class forward-declared in
// local scope
class KernelName2;
kernel_single_task<KernelName2>([=]() { acc.use(); });
#endif

#ifdef LD__
Copy link
Contributor

@kbobrovs kbobrovs Feb 19, 2020

Choose a reason for hiding this comment

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

I suggest to convert these ("expected compilation error") to negative compilation tests together with removal. Or leave it as is if there are no resource now.

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 preserved these negative test cases in the new test, this is mentioned in commit message

// Expected compilation error.
// LD--
// kernel name is a class defined in local scope
class KernelName2a {};
kernel_single_task<KernelName2a>([=]() { acc.use(); });
#endif

// TI--
// an incomplete class forward-declared in a namespace at
// translation unit scope
Expand Down Expand Up @@ -127,16 +111,6 @@ struct MyWrapper {
kernel_single_task<nm1::KernelName3<class KernelName5>>(
[=]() { acc.use(); });

#ifdef TILI
// Expected compilation error
// TILI
// an incomplete template specialization class with incomplete class as
// argument forward-declared locally
class KernelName6;
kernel_single_task<nm1::KernelName3<KernelName6>>(
[=]() { acc.use(); });
#endif

// TDPI
// a defined template specialization class with incomplete class as
// argument forward-declared "in-place"
Expand All @@ -149,40 +123,6 @@ struct MyWrapper {
kernel_single_task<nm1::KernelName8<nm1::nm2::C>>(
[=]() { acc.use(); });

#ifdef TDLI
// TODO unexpected compilation error when host code + integration header
// is compiled TDLI a defined template specialization class with
// incomplete class as argument forward-declared locally
class KernelName6a;
kernel_single_task<nm1::KernelName4<KernelName6a>>(
[=]() { acc.use(); });
#endif

#ifdef TDLD
// Expected compilation error
// TDLD
// a defined template specialization class with a class as argument
// defined locally
class KernelName9 {};
kernel_single_task<nm1::KernelName4<KernelName9>>([=]() { acc.use(); });
#endif

#ifdef TICD
// Expected compilation error
// TICD
// an incomplete template specialization class with a defined class as
// argument declared in the containing class
kernel_single_task<nm1::KernelName3<KN100>>([=]() { acc.use(); });
#endif

#ifdef TICI
// Expected compilation error
// TICI
// an incomplete template specialization class with an incomplete class as
// argument declared in the containing class
kernel_single_task<nm1::KernelName3<KN101>>([=]() { acc.use(); });
#endif

// kernel name type is a templated class, both the top-level class and the
// template argument are declared in the anonymous namespace
kernel_single_task<TmplClassInAnonNS<class ClassInAnonNS>>(
Expand Down
91 changes: 91 additions & 0 deletions clang/test/SemaSYCL/unnamed-kernel.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,91 @@
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -verify %s
// RUN: %clang_cc1 -I %S/Inputs -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -verify %s
#include <sycl.hpp>

#ifdef __SYCL_UNNAMED_LAMBDA__
// expected-no-diagnostics
#endif

namespace namespace1 {
template <typename T>
class KernelName;
}

struct MyWrapper {
private:
class InvalidKernelName0 {};
class InvalidKernelName3 {};
class InvalidKernelName4 {};
class InvalidKernelName5 {};

public:
void test() {
cl::sycl::queue q;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+5 {{kernel needs to have a globally-visible name}}
// expected-note@+2 {{InvalidKernelName1 declared here}}
#endif
class InvalidKernelName1 {};
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidKernelName1>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+5 {{kernel needs to have a globally-visible name}}
// expected-note@+2 {{InvalidKernelName2 declared here}}
#endif
class InvalidKernelName2 {};
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidKernelName2>>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@16 {{InvalidKernelName0 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidKernelName0>([] {});
});

#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@17 {{InvalidKernelName3 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidKernelName3>>([] {});
});

using ValidAlias = MyWrapper;
q.submit([&](cl::sycl::handler &h) {
h.single_task<ValidAlias>([] {});
});

using InvalidAlias = InvalidKernelName4;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@18 {{InvalidKernelName4 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<InvalidAlias>([] {});
});

using InvalidAlias1 = InvalidKernelName5;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+4 {{kernel needs to have a globally-visible name}}
// expected-note@19 {{InvalidKernelName5 declared here}}
#endif
q.submit([&](cl::sycl::handler &h) {
h.single_task<namespace1::KernelName<InvalidAlias1>>([] {});
});
}
};

int main() {
cl::sycl::queue q;
#ifndef __SYCL_UNNAMED_LAMBDA__
// expected-error@+2 {{kernel name is missing}}
#endif
q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); });

return 0;
}