From 4f6ae9129ce152fba3b5d1c19966578d7c6f9d59 Mon Sep 17 00:00:00 2001 From: Srividya Sundaram Date: Wed, 28 Oct 2020 20:51:06 -0700 Subject: [PATCH] [SYCL] Move diagnostic from integration header emission (#2644) This patch moves the diagnostic for the kernel name (kernel name cannot be a type in the "std" namespace) to the point in which the kernel is called, rather than during integration header emission. --- clang/lib/Sema/SemaSYCL.cpp | 59 +++++++++++-------- .../test/CodeGenSYCL/stdtypes_kernel_type.cpp | 57 ------------------ .../template-template-parameter.cpp | 2 +- clang/test/SemaSYCL/stdtypes_kernel_type.cpp | 55 +++++++++++++++++ clang/test/SemaSYCL/unnamed-kernel.cpp | 7 ++- 5 files changed, 96 insertions(+), 84 deletions(-) delete mode 100644 clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp create mode 100644 clang/test/SemaSYCL/stdtypes_kernel_type.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index d08f4bc9ae614..e47d022aac95a 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -2834,10 +2834,15 @@ class SYCLKernelNameTypeVisitor Sema &S; SourceLocation KernelInvocationFuncLoc; using InnerTypeVisitor = TypeVisitor; - using InnerTAVisitor = + using InnerTemplArgVisitor = ConstTemplateArgumentVisitor; bool IsInvalid = false; + void VisitTemplateArgs(ArrayRef Args) { + for (auto &A : Args) + Visit(A); + } + public: SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc) : S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc) {} @@ -2848,15 +2853,19 @@ class SYCLKernelNameTypeVisitor if (T.isNull()) return; const CXXRecordDecl *RD = T->getAsCXXRecordDecl(); - if (!RD) + if (!RD) { + if (T->isNullPtrType()) { + S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named) + << /* kernel name cannot be a type in the std namespace */ 3; + IsInvalid = true; + } return; + } // If KernelNameType has template args visit each template arg via // ConstTemplateArgumentVisitor if (const auto *TSD = dyn_cast(RD)) { - const TemplateArgumentList &Args = TSD->getTemplateArgs(); - for (unsigned I = 0; I < Args.size(); I++) { - Visit(Args[I]); - } + ArrayRef Args = TSD->getTemplateArgs().asArray(); + VisitTemplateArgs(Args); } else { InnerTypeVisitor::Visit(T.getTypePtr()); } @@ -2865,7 +2874,7 @@ class SYCLKernelNameTypeVisitor void Visit(const TemplateArgument &TA) { if (TA.isNull()) return; - InnerTAVisitor::Visit(TA); + InnerTemplArgVisitor::Visit(TA); } void VisitEnumType(const EnumType *T) { @@ -2886,14 +2895,24 @@ class SYCLKernelNameTypeVisitor void VisitTagDecl(const TagDecl *Tag) { bool UnnamedLambdaEnabled = S.getASTContext().getLangOpts().SYCLUnnamedLambda; - if (!Tag->getDeclContext()->isTranslationUnit() && - !isa(Tag->getDeclContext()) && !UnnamedLambdaEnabled) { - const bool KernelNameIsMissing = Tag->getName().empty(); - if (KernelNameIsMissing) { + const DeclContext *DeclCtx = Tag->getDeclContext(); + if (DeclCtx && !UnnamedLambdaEnabled) { + auto *NameSpace = dyn_cast_or_null(DeclCtx); + if (NameSpace && NameSpace->isStdNamespace()) { S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named) - << /* kernel name is missing */ 0; + << /* kernel name cannot be a type in the std namespace */ 3; IsInvalid = true; - } else { + return; + } + if (!DeclCtx->isTranslationUnit() && !isa(DeclCtx)) { + const bool KernelNameIsMissing = Tag->getName().empty(); + if (KernelNameIsMissing) { + S.Diag(KernelInvocationFuncLoc, + diag::err_sycl_kernel_incorrectly_named) + << /* kernel name is missing */ 0; + IsInvalid = true; + return; + } if (Tag->isCompleteDefinition()) { S.Diag(KernelInvocationFuncLoc, diag::err_sycl_kernel_incorrectly_named) @@ -2901,7 +2920,6 @@ class SYCLKernelNameTypeVisitor IsInvalid = true; } else S.Diag(KernelInvocationFuncLoc, diag::warn_sycl_implicit_decl); - S.Diag(Tag->getSourceRange().getBegin(), diag::note_previous_decl) << Tag->getName(); } @@ -2932,6 +2950,10 @@ class SYCLKernelNameTypeVisitor VisitEnumType(ET); } } + + void VisitPackTemplateArgument(const TemplateArgument &TA) { + VisitTemplateArgs(TA.getPackAsArray()); + } }; void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, SourceRange CallLoc, @@ -3337,12 +3359,6 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D, break; } - if (NS->isStdNamespace()) { - Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) - << /* name cannot be a type in the std namespace */ 3; - return; - } - ++NamespaceCnt; const StringRef NSInlinePrefix = NS->isInline() ? "inline " : ""; NSStr.insert( @@ -3426,9 +3442,6 @@ void SYCLIntegrationHeader::emitForwardClassDecls( const CXXRecordDecl *RD = T->getAsCXXRecordDecl(); if (!RD) { - if (T->isNullPtrType()) - Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) - << /* name cannot be a type in the std namespace */ 3; return; } diff --git a/clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp b/clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp deleted file mode 100644 index 475717f93ea85..0000000000000 --- a/clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp +++ /dev/null @@ -1,57 +0,0 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -DCHECK_ERROR -verify %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -triple spir64-unknown-unknown-sycldevice -fsycl-int-header=%t.h %s -// RUN: FileCheck -input-file=%t.h %s -// -// CHECK: #include -// CHECK-NEXT: #include -// -// CHECK: static constexpr -// CHECK-NEXT: const char* const kernel_names[] = { -// CHECK-NEXT: "_ZTSm", -// CHECK-NEXT: "_ZTSl" -// CHECK-NEXT: }; -// -// CHECK: static constexpr -// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { -// CHECK-NEXT: //--- _ZTSm -// CHECK-EMPTY: -// CHECK-NEXT: //--- _ZTSl -// CHECK-EMPTY: -// CHECK-NEXT: }; - -// CHECK: template <> struct KernelInfo { -// CHECK: template <> struct KernelInfo { - -void usage() { -} - -namespace std { -typedef long unsigned int size_t; -typedef long int ptrdiff_t; -typedef decltype(nullptr) nullptr_t; -class T; -class U; -} // namespace std - -template -struct Templated_kernel_name; - -template -__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { - kernelFunc(); -} - -int main() { -#ifdef CHECK_ERROR - kernel_single_task([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}} - kernel_single_task([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}} - kernel_single_task>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}} - kernel_single_task>([=]() {}); // expected-error {{kernel name cannot be a type in the "std" namespace}} -#endif - - // Although in the std namespace, these resolve to builtins such as `int` that are allowed in kernel names - kernel_single_task([=]() {}); - kernel_single_task([=]() {}); - - return 0; -} diff --git a/clang/test/CodeGenSYCL/template-template-parameter.cpp b/clang/test/CodeGenSYCL/template-template-parameter.cpp index 6d5c0491d8941..1a23e30a590b5 100644 --- a/clang/test/CodeGenSYCL/template-template-parameter.cpp +++ b/clang/test/CodeGenSYCL/template-template-parameter.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -sycl-std=2020 %s // RUN: FileCheck -input-file=%t.h %s #include "Inputs/sycl.hpp" diff --git a/clang/test/SemaSYCL/stdtypes_kernel_type.cpp b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp new file mode 100644 index 0000000000000..c4e84e86bea69 --- /dev/null +++ b/clang/test/SemaSYCL/stdtypes_kernel_type.cpp @@ -0,0 +1,55 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -sycl-std=2020 -DCHECK_ERROR -verify %s + +#include "Inputs/sycl.hpp" + +namespace std { +typedef long unsigned int size_t; +typedef long int ptrdiff_t; +typedef decltype(nullptr) nullptr_t; +class T; +class U; +} // namespace std + +template +struct Templated_kernel_name; + +template class TemplParamPack; + +using namespace cl::sycl; +queue q; + +int main() { +#ifdef CHECK_ERROR + // expected-error@Inputs/sycl.hpp:220 5 {{kernel name cannot be a type in the "std" namespace}} + q.submit([&](handler &h) { + // expected-note@+1{{in instantiation of function template specialization}} + h.single_task([=] {}); + }); + q.submit([&](handler &h) { + // expected-note@+1{{in instantiation of function template specialization}} + h.single_task([=] {}); + }); + q.submit([&](handler &h) { + // expected-note@+1{{in instantiation of function template specialization}} + h.single_task>([=] {}); + }); + q.submit([&](handler &h) { + // expected-note@+1{{in instantiation of function template specialization}} + h.single_task>([=] {}); + }); + q.submit([&](handler &cgh) { + // expected-note@+1{{in instantiation of function template specialization}} + cgh.single_task>([]() {}); + }); +#endif + + // Although in the std namespace, these resolve to builtins such as `int` that are allowed in kernel names + q.submit([&](handler &h) { + h.single_task([=] {}); + }); + q.submit([&](handler &h) { + h.single_task([=] {}); + }); + + return 0; +} diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index b90a82d87f933..e8df9198503bc 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -1,5 +1,5 @@ -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsyntax-only -Wno-sycl-2017-compat -verify %s -// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h -fsycl-unnamed-lambda -fsyntax-only -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-unnamed-lambda -fsyntax-only -Wno-sycl-2017-compat -verify %s #include "Inputs/sycl.hpp" #ifdef __SYCL_UNNAMED_LAMBDA__ @@ -70,7 +70,8 @@ struct MyWrapper { }); #ifndef __SYCL_UNNAMED_LAMBDA__ - // expected-error@+3 {{kernel name cannot be a type in the "std" namespace}} + // expected-error@Inputs/sycl.hpp:220 {{kernel name cannot be a type in the "std" namespace}} + // expected-note@+3{{in instantiation of function template specialization}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {});