From 5c4a84bbe2e9387e29297130daf53e613b883bea Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Wed, 5 Jun 2019 08:44:44 -0700 Subject: [PATCH 1/2] [SYCL] Support unnamed lambda kernels Add support for kernels without kernel name. Requires compiling with -DUNNAMED_LAMBDA_EXT. Uses new __unique_stable_name compiler built-in. Signed-off-by: Roland Schulz Signed-off-by: Alexey Bader --- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/include/clang/Sema/Sema.h | 6 +- clang/lib/Sema/SemaSYCL.cpp | 26 ++++--- clang/test/CodeGenSYCL/integration_header.cpp | 1 - clang/test/CodeGenSYCL/wrapped-accessor.cpp | 2 - clang/test/Misc/warning-flags.c | 3 +- sycl/include/CL/sycl/detail/kernel_desc.hpp | 24 +++++++ sycl/test/regression/kernel_name_class.cpp | 2 + sycl/test/regression/kernel_unnamed.cpp | 67 +++++++++++++++++++ 9 files changed, 118 insertions(+), 15 deletions(-) create mode 100644 sycl/test/regression/kernel_unnamed.cpp diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index f950c03ec9dbc..cdc7a5eb819d2 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9768,7 +9768,7 @@ 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< +def warn_sycl_kernel_name_class_not_top_level : Warning< "kernel name class and its template argument classes' declarations can only " "nest in a namespace: %0">; def err_sycl_restrict : Error< diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 6efeb1c0959a6..65deceb8e2a15 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -317,7 +317,8 @@ 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); + void startKernel(StringRef KernelName, QualType KernelNameType, + StringRef KernelStableName); /// Adds a kernel parameter descriptor to current kernel invocation /// descriptor. @@ -352,6 +353,9 @@ class SYCLIntegrationHeader { /// Kernel name type. QualType NameType; + /// Kernel name with stable lamba name mangling + std::string StableName; + /// Descriptor of kernel actual parameters. SmallVector Params; diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 73fdb7dd5c2cd..4124a354a64c5 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -797,7 +797,9 @@ static void populateIntHeader(SYCLIntegrationHeader &H, const StringRef Name, ASTContext &Ctx = KernelObjTy->getASTContext(); const ASTRecordLayout &Layout = Ctx.getASTRecordLayout(KernelObjTy); - H.startKernel(Name, NameType); + const std::string StableName = PredefinedExpr::ComputeName( + Ctx, PredefinedExpr::UniqueStableNameExpr, NameType); + H.startKernel(Name, NameType, StableName); auto populateHeaderForAccessor = [&](const QualType &ArgTy, uint64_t Offset) { // The parameter is a SYCL accessor object. @@ -1112,7 +1114,7 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) { // 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); + diag::warn_sycl_kernel_name_class_not_top_level); } } break; @@ -1238,12 +1240,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "#include \n"; O << "\n"; + O << "#ifndef UNNAMED_LAMBDA_EXT\n"; O << "// Forward declarations of templated kernel function types:\n"; llvm::SmallPtrSet Printed; for (const KernelDesc &K : KernelDescs) { emitForwardClassDecls(O, K.NameType, Printed); } + O << "#endif\n"; O << "\n"; O << "namespace cl {\n"; @@ -1305,19 +1309,21 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { } O << "};\n\n"; - O << "// Specializations of this template class encompasses information\n"; - O << "// about a kernel. The kernel is identified by the template\n"; - O << "// parameter type.\n"; - O << "template struct KernelInfo;\n"; - O << "\n"; - O << "// Specializations of KernelInfo for kernel function types:\n"; CurStart = 0; for (const KernelDesc &K : KernelDescs) { const size_t N = K.Params.size(); + O << "#ifdef UNNAMED_LAMBDA_EXT\n"; + O << "template <> struct KernelInfoData<"; + O << "'" << K.StableName.front(); + for (char c : StringRef(K.StableName).substr(1)) + O << "', '" << c; + O << "'> {\n"; + O << "#else\n"; O << "template <> struct KernelInfo<" << eraseAnonNamespace(K.NameType.getAsString()) << "> {\n"; + O << "#endif\n"; O << " DLL_LOCAL\n"; O << " static constexpr const char* getName() { return \"" << K.Name << "\"; }\n"; @@ -1355,10 +1361,12 @@ bool SYCLIntegrationHeader::emit(const StringRef &IntHeaderName) { } void SYCLIntegrationHeader::startKernel(StringRef KernelName, - QualType KernelNameType) { + QualType KernelNameType, + StringRef KernelStableName) { KernelDescs.resize(KernelDescs.size() + 1); KernelDescs.back().Name = KernelName; KernelDescs.back().NameType = KernelNameType; + KernelDescs.back().StableName = KernelStableName; } void SYCLIntegrationHeader::addParamDesc(kernel_param_kind_t Kind, int Info, diff --git a/clang/test/CodeGenSYCL/integration_header.cpp b/clang/test/CodeGenSYCL/integration_header.cpp index e0ef5f6cffda2..92ce63de9fb6b 100644 --- a/clang/test/CodeGenSYCL/integration_header.cpp +++ b/clang/test/CodeGenSYCL/integration_header.cpp @@ -47,7 +47,6 @@ // CHECK-EMPTY: // CHECK-NEXT: }; // -// CHECK: template struct KernelInfo; // CHECK: template <> struct KernelInfo { // CHECK: template <> struct KernelInfo<::second_namespace::second_kernel> { // CHECK: template <> struct KernelInfo<::third_kernel<1, int, ::point >> { diff --git a/clang/test/CodeGenSYCL/wrapped-accessor.cpp b/clang/test/CodeGenSYCL/wrapped-accessor.cpp index 2b1d8e176013a..f286af7398b93 100644 --- a/clang/test/CodeGenSYCL/wrapped-accessor.cpp +++ b/clang/test/CodeGenSYCL/wrapped-accessor.cpp @@ -27,8 +27,6 @@ // CHECK-NEXT: 0 // _ZTSZ4mainE14wrapped_access // CHECK-NEXT: }; -// CHECK: template struct KernelInfo; - // CHECK: template <> struct KernelInfo { #include diff --git a/clang/test/Misc/warning-flags.c b/clang/test/Misc/warning-flags.c index 81d332cacd419..abb60f58ea386 100644 --- a/clang/test/Misc/warning-flags.c +++ b/clang/test/Misc/warning-flags.c @@ -18,7 +18,7 @@ This test serves two purposes: The list of warnings below should NEVER grow. It should gradually shrink to 0. -CHECK: Warnings without flags (74): +CHECK: Warnings without flags (75): CHECK-NEXT: ext_excess_initializers CHECK-NEXT: ext_excess_initializers_in_char_array_initializer CHECK-NEXT: ext_expected_semi_decl_list @@ -84,6 +84,7 @@ CHECK-NEXT: warn_property_getter_owning_mismatch CHECK-NEXT: warn_register_objc_catch_parm CHECK-NEXT: warn_related_result_type_compatibility_class CHECK-NEXT: warn_related_result_type_compatibility_protocol +CHECK-NEXT: warn_sycl_kernel_name_class_not_top_level CHECK-NEXT: warn_template_export_unsupported CHECK-NEXT: warn_template_spec_extra_headers CHECK-NEXT: warn_tentative_incomplete_array diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index ebf71ac8561c7..a7aa9104e9789 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -51,6 +51,7 @@ struct kernel_param_desc_t { int offset; }; +#ifndef UNNAMED_LAMBDA_EXT template struct KernelInfo { static constexpr unsigned getNumParams() { return 0; } static const kernel_param_desc_t &getParamDesc(int Idx) { @@ -59,6 +60,29 @@ template struct KernelInfo { } static constexpr const char *getName() { return ""; } }; +#else +template struct KernelInfoData; // Should this have dummy impl? + +// C++14 like index_sequence and make_index_sequence +// not needed C++14 members (value_type, size) not implemented +template struct integer_sequence {}; +template using index_sequence = integer_sequence; +template +using make_index_sequence = __make_integer_seq; + +template struct KernelInfoImpl { +private: + static constexpr auto n = __unique_stable_name(T); + template + static KernelInfoData impl(index_sequence) { + return {}; + } + +public: + using type = decltype(impl(make_index_sequence<__builtin_strlen(n)>{})); +}; +template using KernelInfo = typename KernelInfoImpl::type; +#endif } // namespace detail } // namespace sycl diff --git a/sycl/test/regression/kernel_name_class.cpp b/sycl/test/regression/kernel_name_class.cpp index b38e2b2b2574c..9eb5d2de56211 100644 --- a/sycl/test/regression/kernel_name_class.cpp +++ b/sycl/test/regression/kernel_name_class.cpp @@ -3,6 +3,8 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out +// RUN: %clangxx -fsycl %s -o %t.ext.out -lOpenCL -DUNNAMED_LAMBDA_EXT +// RUN: %CPU_RUN_PLACEHOLDER %t.ext.out //==-- kernel_name_class.cpp - SYCL kernel naming variants test ------------==// // diff --git a/sycl/test/regression/kernel_unnamed.cpp b/sycl/test/regression/kernel_unnamed.cpp new file mode 100644 index 0000000000000..4568247216a23 --- /dev/null +++ b/sycl/test/regression/kernel_unnamed.cpp @@ -0,0 +1,67 @@ +// RUN: %clangxx -fsycl %s -o %t.out -lOpenCL -DUNNAMED_LAMBDA_EXT +// RUN: env SYCL_DEVICE_TYPE=HOST %t.out +// RUN: %CPU_RUN_PLACEHOLDER %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out +// RUN: %ACC_RUN_PLACEHOLDER %t.out + +//==-- kernel_unnamed.cpp - SYCL kernel naming variants test ------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include + +#define GOLD 10 +static int NumTestCases = 0; + +template +void foo(cl::sycl::queue &deviceQueue, cl::sycl::buffer &buf, F f) { + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { acc[0] = f(acc[0], GOLD); }); + }); +} + +namespace nm { +struct Wrapper { + + int test() { + int arr[] = {0}; + { + // Simple test + cl::sycl::queue deviceQueue; + cl::sycl::buffer buf(arr, 1); + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { acc[0] += GOLD; }); + }); + ++NumTestCases; + +// Test lambdas with different ordinal because of macro expansion +#ifdef __SYCL_DEVICE_ONLY__ + [] {}(); +#endif + deviceQueue.submit([&](cl::sycl::handler &cgh) { + auto acc = buf.get_access(cgh); + cgh.single_task([=]() { acc[0] += GOLD; }); + }); + ++NumTestCases; + + // Test lambda passed to function + foo(deviceQueue, buf, [](int a, int b) { return a + b; }); + ++NumTestCases; + } + return arr[0]; + } +}; +} // namespace nm + +int main() { + nm::Wrapper w; + int res = w.test(); + assert (res == GOLD * NumTestCases && "Wrong result"); +} From 372222daa765bb2a61be8b70dce2efa18e711014 Mon Sep 17 00:00:00 2001 From: Roland Schulz Date: Thu, 25 Jul 2019 21:30:41 -0700 Subject: [PATCH 2/2] [SYCL] Add -fsycl-unnamed-lambda flag User activates the extension with the new flag rather than previously with -DUNNAMED_LAMBDA_EXT. Fixes that the err_sycl_kernel_name_class_not_top_level is an error without the extension and not a warning with it. Only the version of the integration header needed is generated (rather than both version with #ifdef). Defines __SYCL_UNNAMED_LAMBDA__ when extension is active. Signed-off-by: Roland Schulz --- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/include/clang/Basic/LangOptions.def | 1 + clang/include/clang/Driver/Options.td | 2 + clang/include/clang/Sema/Sema.h | 9 ++-- clang/lib/Driver/ToolChains/Clang.cpp | 2 + clang/lib/Frontend/CompilerInvocation.cpp | 1 + clang/lib/Frontend/InitPreprocessor.cpp | 2 + clang/lib/Sema/SemaSYCL.cpp | 41 ++++++++++--------- clang/test/Misc/warning-flags.c | 3 +- sycl/include/CL/sycl/detail/kernel_desc.hpp | 4 +- sycl/test/regression/kernel_name_class.cpp | 2 +- sycl/test/regression/kernel_unnamed.cpp | 2 +- 12 files changed, 41 insertions(+), 30 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index cdc7a5eb819d2..f950c03ec9dbc 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -9768,7 +9768,7 @@ 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 warn_sycl_kernel_name_class_not_top_level : Warning< +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_restrict : Error< diff --git a/clang/include/clang/Basic/LangOptions.def b/clang/include/clang/Basic/LangOptions.def index cdda3145c5b4e..671bc061da9ac 100644 --- a/clang/include/clang/Basic/LangOptions.def +++ b/clang/include/clang/Basic/LangOptions.def @@ -225,6 +225,7 @@ LANGOPT(GPURelocatableDeviceCode, 1, 0, "generate relocatable device code") LANGOPT(SYCLIsDevice , 1, 0, "Generate code for SYCL device") LANGOPT(SYCLIsHost , 1, 0, "SYCL host compilation") LANGOPT(SYCLAllowFuncPtr , 1, 0, "Allow function pointers in SYCL device code") +LANGOPT(SYCLUnnamedLambda , 1, 0, "Allow unnamed lambda SYCL kernels") LANGOPT(SizedDeallocation , 1, 0, "sized deallocation") LANGOPT(AlignedAllocation , 1, 0, "aligned allocation") diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index 4c18f8cbb81a9..e999d33d549c6 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -1739,6 +1739,8 @@ def fno_sycl_use_bitcode : Flag<["-"], "fno-sycl-use-bitcode">, Flags<[CC1Option]>, HelpText<"Use SPIR-V instead of LLVM bitcode in fat objects">; def fsycl_link : Flag<["-"], "fsycl-link">, Flags<[CC1Option]>, HelpText<"Generate partially linked device object to be used with the host link">; +def fsycl_unnamed_lambda : Flag<["-"], "fsycl-unnamed-lambda">, + Flags<[CC1Option]>, HelpText<"Allow unnamed SYCL lambda kernels">; def fsyntax_only : Flag<["-"], "fsyntax-only">, Flags<[DriverOption,CoreOption,CC1Option]>, Group; def ftabstop_EQ : Joined<["-"], "ftabstop=">, Group; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 65deceb8e2a15..9c13b62c8dd11 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -306,7 +306,7 @@ class SYCLIntegrationHeader { }; public: - SYCLIntegrationHeader(DiagnosticsEngine &Diag); + SYCLIntegrationHeader(DiagnosticsEngine &Diag, bool UnnamedLambdaSupport); /// Emits contents of the header into given stream. void emit(raw_ostream &Out); @@ -353,7 +353,7 @@ class SYCLIntegrationHeader { /// Kernel name type. QualType NameType; - /// Kernel name with stable lamba name mangling + /// Kernel name with stable lambda name mangling std::string StableName; /// Descriptor of kernel actual parameters. @@ -391,6 +391,9 @@ class SYCLIntegrationHeader { /// Used for emitting diagnostics. DiagnosticsEngine &Diag; + + /// Whether header is generated with unnamed lambda support + bool UnnamedLambdaSupport; }; /// Keeps track of expected type during expression parsing. The type is tied to @@ -11408,7 +11411,7 @@ class Sema { SYCLIntegrationHeader &getSyclIntegrationHeader() { if (SyclIntHeader == nullptr) SyclIntHeader = llvm::make_unique( - getDiagnostics()); + getDiagnostics(), getLangOpts().SYCLUnnamedLambda); return *SyclIntHeader.get(); } diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index be7b9aac9d952..cb0f16d4d4896 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5363,6 +5363,8 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, HeaderOpt += Output.getFilename(); CmdArgs.push_back(Args.MakeArgString(HeaderOpt)); } + if (Args.hasArg(options::OPT_fsycl_unnamed_lambda)) + CmdArgs.push_back("-fsycl-unnamed-lambda"); } // OpenMP offloading device jobs take the argument -fopenmp-host-ir-file-path diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index 151b96ddf9109..35a402f7473fc 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -2973,6 +2973,7 @@ static void ParseLangArgs(LangOptions &Opts, ArgList &Args, InputKind IK, Opts.SYCLIsHost = Args.hasArg(options::OPT_fsycl_is_host); Opts.SYCLAllowFuncPtr = Args.hasFlag(options::OPT_fsycl_allow_func_ptr, options::OPT_fno_sycl_allow_func_ptr, false); + Opts.SYCLUnnamedLambda = Args.hasArg(options::OPT_fsycl_unnamed_lambda); // Set CUDA mode for OpenMP target NVPTX if specified in options Opts.OpenMPCUDAMode = Opts.OpenMPIsDevice && T.isNVPTX() && diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp index 6a72b40f0bef8..4d4bd13bb85a6 100644 --- a/clang/lib/Frontend/InitPreprocessor.cpp +++ b/clang/lib/Frontend/InitPreprocessor.cpp @@ -1065,6 +1065,8 @@ static void InitializePredefinedMacros(const TargetInfo &TI, if (!getenv("DISABLE_INFER_AS")) Builder.defineMacro("__SYCL_ENABLE_INFER_AS__", "1"); } + if (LangOpts.SYCLUnnamedLambda) + Builder.defineMacro("__SYCL_UNNAMED_LAMBDA__", "1"); // OpenCL definitions. if (LangOpts.OpenCL) { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 4124a354a64c5..c2e83dd85c3ed 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1110,11 +1110,11 @@ void SYCLIntegrationHeader::emitFwdDecl(raw_ostream &O, const Decl *D) { ? cast(D)->getTemplatedDecl() : dyn_cast(D); - if (TD && TD->isCompleteDefinition()) { + if (TD && TD->isCompleteDefinition() && !UnnamedLambdaSupport) { // defined class constituting the kernel name is not globally // accessible - contradicts the spec Diag.Report(D->getSourceRange().getBegin(), - diag::warn_sycl_kernel_name_class_not_top_level); + diag::err_sycl_kernel_name_class_not_top_level); } } break; @@ -1240,14 +1240,14 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << "#include \n"; O << "\n"; - O << "#ifndef UNNAMED_LAMBDA_EXT\n"; - O << "// Forward declarations of templated kernel function types:\n"; + if (!UnnamedLambdaSupport) { + O << "// Forward declarations of templated kernel function types:\n"; - llvm::SmallPtrSet Printed; - for (const KernelDesc &K : KernelDescs) { - emitForwardClassDecls(O, K.NameType, Printed); + llvm::SmallPtrSet Printed; + for (const KernelDesc &K : KernelDescs) { + emitForwardClassDecls(O, K.NameType, Printed); + } } - O << "#endif\n"; O << "\n"; O << "namespace cl {\n"; @@ -1314,16 +1314,16 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { for (const KernelDesc &K : KernelDescs) { const size_t N = K.Params.size(); - O << "#ifdef UNNAMED_LAMBDA_EXT\n"; - O << "template <> struct KernelInfoData<"; - O << "'" << K.StableName.front(); - for (char c : StringRef(K.StableName).substr(1)) - O << "', '" << c; - O << "'> {\n"; - O << "#else\n"; - O << "template <> struct KernelInfo<" - << eraseAnonNamespace(K.NameType.getAsString()) << "> {\n"; - O << "#endif\n"; + if (UnnamedLambdaSupport) { + O << "template <> struct KernelInfoData<"; + O << "'" << K.StableName.front(); + for (char c : StringRef(K.StableName).substr(1)) + O << "', '" << c; + O << "'> {\n"; + } else { + O << "template <> struct KernelInfo<" + << eraseAnonNamespace(K.NameType.getAsString()) << "> {\n"; + } O << " DLL_LOCAL\n"; O << " static constexpr const char* getName() { return \"" << K.Name << "\"; }\n"; @@ -1384,8 +1384,9 @@ void SYCLIntegrationHeader::endKernel() { // nop for now } -SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag) - : Diag(_Diag) {} +SYCLIntegrationHeader::SYCLIntegrationHeader(DiagnosticsEngine &_Diag, + bool _UnnamedLambdaSupport) + : Diag(_Diag), UnnamedLambdaSupport(_UnnamedLambdaSupport) {} bool Util::isSyclAccessorType(const QualType &Ty) { static std::array Scopes = { diff --git a/clang/test/Misc/warning-flags.c b/clang/test/Misc/warning-flags.c index abb60f58ea386..81d332cacd419 100644 --- a/clang/test/Misc/warning-flags.c +++ b/clang/test/Misc/warning-flags.c @@ -18,7 +18,7 @@ This test serves two purposes: The list of warnings below should NEVER grow. It should gradually shrink to 0. -CHECK: Warnings without flags (75): +CHECK: Warnings without flags (74): CHECK-NEXT: ext_excess_initializers CHECK-NEXT: ext_excess_initializers_in_char_array_initializer CHECK-NEXT: ext_expected_semi_decl_list @@ -84,7 +84,6 @@ CHECK-NEXT: warn_property_getter_owning_mismatch CHECK-NEXT: warn_register_objc_catch_parm CHECK-NEXT: warn_related_result_type_compatibility_class CHECK-NEXT: warn_related_result_type_compatibility_protocol -CHECK-NEXT: warn_sycl_kernel_name_class_not_top_level CHECK-NEXT: warn_template_export_unsupported CHECK-NEXT: warn_template_spec_extra_headers CHECK-NEXT: warn_tentative_incomplete_array diff --git a/sycl/include/CL/sycl/detail/kernel_desc.hpp b/sycl/include/CL/sycl/detail/kernel_desc.hpp index a7aa9104e9789..5ff253ad8e27b 100644 --- a/sycl/include/CL/sycl/detail/kernel_desc.hpp +++ b/sycl/include/CL/sycl/detail/kernel_desc.hpp @@ -51,7 +51,7 @@ struct kernel_param_desc_t { int offset; }; -#ifndef UNNAMED_LAMBDA_EXT +#ifndef __SYCL_UNNAMED_LAMBDA__ template struct KernelInfo { static constexpr unsigned getNumParams() { return 0; } static const kernel_param_desc_t &getParamDesc(int Idx) { @@ -82,7 +82,7 @@ template struct KernelInfoImpl { using type = decltype(impl(make_index_sequence<__builtin_strlen(n)>{})); }; template using KernelInfo = typename KernelInfoImpl::type; -#endif +#endif //__SYCL_UNNAMED_LAMBDA__ } // namespace detail } // namespace sycl diff --git a/sycl/test/regression/kernel_name_class.cpp b/sycl/test/regression/kernel_name_class.cpp index 9eb5d2de56211..9f3959aa2912a 100644 --- a/sycl/test/regression/kernel_name_class.cpp +++ b/sycl/test/regression/kernel_name_class.cpp @@ -3,7 +3,7 @@ // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out // RUN: %ACC_RUN_PLACEHOLDER %t.out -// RUN: %clangxx -fsycl %s -o %t.ext.out -lOpenCL -DUNNAMED_LAMBDA_EXT +// RUN: %clangxx -fsycl %s -o %t.ext.out -lOpenCL -fsycl-unnamed-lambda // RUN: %CPU_RUN_PLACEHOLDER %t.ext.out //==-- kernel_name_class.cpp - SYCL kernel naming variants test ------------==// diff --git a/sycl/test/regression/kernel_unnamed.cpp b/sycl/test/regression/kernel_unnamed.cpp index 4568247216a23..9b2bb4c870c58 100644 --- a/sycl/test/regression/kernel_unnamed.cpp +++ b/sycl/test/regression/kernel_unnamed.cpp @@ -1,4 +1,4 @@ -// RUN: %clangxx -fsycl %s -o %t.out -lOpenCL -DUNNAMED_LAMBDA_EXT +// RUN: %clangxx -fsycl %s -o %t.out -lOpenCL -fsycl-unnamed-lambda // RUN: env SYCL_DEVICE_TYPE=HOST %t.out // RUN: %CPU_RUN_PLACEHOLDER %t.out // RUN: %GPU_RUN_PLACEHOLDER %t.out