diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b3ecda585cf3e..1c80f04356bea 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1649,30 +1649,27 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D, auto *NS = dyn_cast_or_null(DC); if (!NS) { - if (!DC->isTranslationUnit()) { - const TagDecl *TD = isa(D) - ? cast(D)->getTemplatedDecl() - : dyn_cast(D); - - if (TD && !UnnamedLambdaSupport) { - // defined class constituting the kernel name is not globally - // accessible - contradicts the spec - const bool KernelNameIsMissing = TD->getName().empty(); - if (KernelNameIsMissing) { + const TagDecl *TD = isa(D) + ? cast(D)->getTemplatedDecl() + : dyn_cast(D); + if (!TD) + break; + + const bool KernelNameIsMissing = TD->getName().empty(); + if (KernelNameIsMissing) + Diag.Report(KernelLocation, diag::err_sycl_kernel_incorrectly_named) + << /* kernel name is missing */ 0; + else if (!DC->isTranslationUnit()) { + // defined class constituting the kernel name is not globally + // accessible - contradicts the spec + if (!UnnamedLambdaSupport) { + if (TD->isCompleteDefinition()) 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 { - if (TD->isCompleteDefinition()) - Diag.Report(KernelLocation, - diag::err_sycl_kernel_incorrectly_named) - << /* kernel name is not globally-visible */ 1; - else - Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl); - Diag.Report(D->getSourceRange().getBegin(), - diag::note_previous_decl) - << TD->getName(); - } + << /* kernel name is not globally-visible */ 1; + else + Diag.Report(KernelLocation, diag::warn_sycl_implicit_decl); + Diag.Report(D->getSourceRange().getBegin(), diag::note_previous_decl) + << TD->getName(); } } break; @@ -1826,8 +1823,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "// This is auto-generated SYCL integration header.\n"; O << "\n"; + O << "#include \n"; O << "#include \n"; O << "#include \n"; + O << "using nullptr_t = std::nullptr_t;\n"; O << "\n"; diff --git a/clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp b/clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp new file mode 100644 index 0000000000000..aa2f5efe4ded0 --- /dev/null +++ b/clang/test/CodeGenSYCL/stdtypes_kernel_type.cpp @@ -0,0 +1,63 @@ +// RUN: %clang_cc1 -fsycl -fsycl-is-device -fsycl-int-header=%t.h %s +// RUN: FileCheck -input-file=%t.h %s +// +// CHECK: #include +// CHECK-NEXT: #include +// CHECK-NEXT: #include +// CHECK-NEXT: using nullptr_t = std::nullptr_t; +// +// CHECK: static constexpr +// CHECK-NEXT: const char* const kernel_names[] = { +// CHECK-NEXT: "_ZTSDn" +// CHECK-NEXT: "_ZTSSt4byte" +// CHECK-NEXT: "_ZTSm", +// CHECK-NEXT: "_ZTSl" +// CHECK-NEXT: }; +// +// CHECK: static constexpr +// CHECK-NEXT: const kernel_param_desc_t kernel_signatures[] = { +// CHECK-NEXT: //--- _ZTSDn +// CHECK-EMPTY: +// CHECK-NEXT: //--- _ZTSSt4byte +// CHECK-EMPTY: +// CHECK-NEXT: //--- _ZTSm +// CHECK-EMPTY: +// CHECK-NEXT: //--- _ZTSl +// CHECK-EMPTY: +// CHECK-NEXT: }; +// +// CHECK: static constexpr +// CHECK-NEXT: const unsigned kernel_signature_start[] = { +// CHECK-NEXT: 0, // _ZTSDn +// CHECK-NEXT: 1, // _ZTSSt4byte +// CHECK-NEXT: 2, // _ZTSm +// CHECK-NEXT: 3 // _ZTSl +// CHECK-NEXT: }; + +// CHECK: template <> struct KernelInfo { +// CHECK: template <> struct KernelInfo<::std::byte> { +// 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; +enum class byte : unsigned char {}; +} // namespace std + +template +__attribute__((sycl_kernel)) void kernel_single_task(Func kernelFunc) { + kernelFunc(); +} + +int main() { + kernel_single_task([]() { usage(); }); + kernel_single_task([=]() {}); + kernel_single_task([=]() {}); + kernel_single_task([=]() {}); + return 0; +} diff --git a/clang/test/SemaSYCL/unnamed-kernel.cpp b/clang/test/SemaSYCL/unnamed-kernel.cpp index 33e5656ac8220..d36c4d13cdd8b 100644 --- a/clang/test/SemaSYCL/unnamed-kernel.cpp +++ b/clang/test/SemaSYCL/unnamed-kernel.cpp @@ -1,16 +1,22 @@ // 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 #ifdef __SYCL_UNNAMED_LAMBDA__ // expected-no-diagnostics #endif +#include + namespace namespace1 { template class KernelName; } +namespace std { +typedef struct { +} max_align_t; +} // namespace std + struct MyWrapper { private: class InvalidKernelName0 {}; @@ -41,7 +47,7 @@ struct MyWrapper { #ifndef __SYCL_UNNAMED_LAMBDA__ // expected-error@+4 {{kernel needs to have a globally-visible name}} - // expected-note@16 {{InvalidKernelName0 declared here}} + // expected-note@21 {{InvalidKernelName0 declared here}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); @@ -49,7 +55,7 @@ struct MyWrapper { #ifndef __SYCL_UNNAMED_LAMBDA__ // expected-error@+4 {{kernel needs to have a globally-visible name}} - // expected-note@17 {{InvalidKernelName3 declared here}} + // expected-note@22 {{InvalidKernelName3 declared here}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task>([] {}); @@ -60,10 +66,17 @@ struct MyWrapper { h.single_task([] {}); }); +#ifndef __SYCL_UNNAMED_LAMBDA__ + // expected-error@+3 {{kernel name is missing}} +#endif + q.submit([&](cl::sycl::handler &h) { + h.single_task([] {}); + }); + using InvalidAlias = InvalidKernelName4; #ifndef __SYCL_UNNAMED_LAMBDA__ // expected-error@+4 {{kernel needs to have a globally-visible name}} - // expected-note@18 {{InvalidKernelName4 declared here}} + // expected-note@23 {{InvalidKernelName4 declared here}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task([] {}); @@ -72,7 +85,7 @@ struct MyWrapper { using InvalidAlias1 = InvalidKernelName5; #ifndef __SYCL_UNNAMED_LAMBDA__ // expected-error@+4 {{kernel needs to have a globally-visible name}} - // expected-note@19 {{InvalidKernelName5 declared here}} + // expected-note@24 {{InvalidKernelName5 declared here}} #endif q.submit([&](cl::sycl::handler &h) { h.single_task>([] {});