Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
28 commits
Select commit Hold shift + click to select a range
27597e2
Removed fsycl-early-opts and reimplemented fno-sycl-early-opts
Nov 28, 2022
8546bec
Fixed faulty flag assignment
Dec 2, 2022
78c9589
Reverted change to sycl-early-opts flag
Dec 6, 2022
d4ae056
Reverted two tests with unnecessary changes
Dec 6, 2022
dd3445e
Fixed the flag definition and tidied up the logic
Dec 7, 2022
2416d7d
clang-format'd BackendUtil changes
Dec 8, 2022
63c4430
Merge branch 'sycl' into alamzeds/fsycl-early-optimizations-flag
Dec 8, 2022
9c8bbb5
Resolved failing clang-format issues
Dec 8, 2022
bc085fd
Changed sycl-early-opts flag back to prev definition w/ change
Dec 9, 2022
95b0d3c
DisableLLVMPasses flag now handled by marshalling infrastructure
Dec 13, 2022
c69f54f
Optimization pipeline refactor to honour flags and fix logic
Dec 13, 2022
b13fa37
Update to group-local-memory test to honour disable-llvm-passes
Dec 13, 2022
af1324b
Rolled back uses_aspect test now it adheres to -disable-llvm-passes
Dec 13, 2022
5f45075
sub-group-size test change with updated flags
Dec 14, 2022
63270a6
Fixed functional pass invocation logic
Dec 14, 2022
7c36bb7
Converted inlining test to check nodes in AST
Dec 16, 2022
aac0dde
Reverted device_has test to respect disable_llvm_passes flag
Dec 19, 2022
b7586ff
Merge branch 'sycl' into alamzeds/fsycl-early-optimizations-flag
Dec 19, 2022
f8c9ab1
Ran clang-format over device_has sycl/test test
Dec 19, 2022
c46a0db
Refactored pipeline building logic after removing first SYCLPropagate…
Dec 20, 2022
639ae36
Merge branch 'sycl' into alamzeds/fsycl-early-optimizations-flag
Dec 20, 2022
5d39281
Tidied up and formatted pipeline if logic
Dec 20, 2022
2c90ec7
Let DisableSYCLEarlyOpts codegen opt be set by marshalling
Jan 3, 2023
786f8c0
Refactor to consolidate logic and clean up code paths
Jan 3, 2023
34c9d24
Merge branch 'sycl' into alamzeds/fsycl-early-optimizations-flag
Jan 3, 2023
c4c5274
Formatted changes
Jan 3, 2023
c2adfd0
Merge branch 'sycl' into alamzeds/fsycl-early-optimizations-flag
Jan 4, 2023
9871f46
Updated force inline kernel lambda test
Jan 5, 2023
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
3 changes: 3 additions & 0 deletions clang/include/clang/Basic/CodeGenOptions.def
Original file line number Diff line number Diff line change
Expand Up @@ -513,6 +513,9 @@ CODEGENOPT(OpaquePointers, 1, 0)
/// non-deleting destructors. (No effect on Microsoft ABI.)
CODEGENOPT(CtorDtorReturnThis, 1, 0)

/// Whether to disable the standard optimization pipeline for the SYCL device compiler.
CODEGENOPT(DisableSYCLEarlyOpts, 1, 0)

#undef CODEGENOPT
#undef ENUM_CODEGENOPT
#undef VALUE_CODEGENOPT
3 changes: 2 additions & 1 deletion clang/include/clang/Driver/Options.td
Original file line number Diff line number Diff line change
Expand Up @@ -5110,7 +5110,8 @@ def : Flag<["-"], "fno-sycl-explicit-simd">,
Flags<[CoreOption, Deprecated]>,
Group<clang_ignored_legacy_options_Group>,
HelpText<"Disable SYCL explicit SIMD extension. (deprecated)">;
defm sycl_early_optimizations : OptOutCC1FFlag<"sycl-early-optimizations", "Enable", "Disable", " standard optimization pipeline for SYCL device compiler", [CoreOption]>;
defm sycl_early_optimizations : OptOutCC1FFlag<"sycl-early-optimizations", "Enable", "Disable", " standard optimization pipeline for SYCL device compiler", [CoreOption]>,
MarshallingInfoFlag<CodeGenOpts<"DisableSYCLEarlyOpts">>;
def fsycl_dead_args_optimization : Flag<["-"], "fsycl-dead-args-optimization">,
Group<sycl_Group>, Flags<[NoArgumentUnused, CoreOption]>, HelpText<"Enables "
"elimination of DPC++ dead kernel arguments">;
Expand Down
51 changes: 22 additions & 29 deletions clang/lib/CodeGen/BackendUtil.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -916,11 +916,6 @@ void EmitAssemblyHelper::RunOptimizationPipeline(

ModulePassManager MPM;

// FIXME: Change this when -fno-sycl-early-optimizations is not tied to
// -disable-llvm-passes.
if (CodeGenOpts.DisableLLVMPasses && LangOpts.SYCLIsDevice)
MPM.addPass(SYCLPropagateAspectsUsagePass());

if (!CodeGenOpts.DisableLLVMPasses) {
// Map our optimization levels into one of the distinct levels used to
// configure the pipeline.
Expand Down Expand Up @@ -1021,7 +1016,10 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
MPM.addPass(InstrProfiling(*Options, false));
});

if (CodeGenOpts.OptimizationLevel == 0) {
if (CodeGenOpts.DisableSYCLEarlyOpts) {
MPM =
PB.buildO0DefaultPipeline(OptimizationLevel::O0, IsLTO || IsThinLTO);
} else if (CodeGenOpts.OptimizationLevel == 0) {
MPM = PB.buildO0DefaultPipeline(Level, IsLTO || IsThinLTO);
} else if (IsThinLTO) {
MPM = PB.buildThinLTOPreLinkDefaultPipeline(Level);
Expand All @@ -1035,31 +1033,26 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
MPM.addPass(createModuleToFunctionPassAdaptor(MemProfilerPass()));
MPM.addPass(ModuleMemProfilerPass());
}
}
if (LangOpts.SYCLIsDevice) {
MPM.addPass(SYCLMutatePrintfAddrspacePass());
if (!CodeGenOpts.DisableLLVMPasses && LangOpts.EnableDAEInSpirKernels)
MPM.addPass(DeadArgumentEliminationSYCLPass());
}

// Add SPIRITTAnnotations pass to the pass manager if
// -fsycl-instrument-device-code option was passed. This option can be used
// only with spir triple.
if (LangOpts.SYCLIsDevice && CodeGenOpts.SPIRITTAnnotations) {
assert(TargetTriple.isSPIR() &&
"ITT annotations can only be added to a module with spir target");
MPM.addPass(SPIRITTAnnotationsPass());
}
if (LangOpts.SYCLIsDevice) {
MPM.addPass(SYCLMutatePrintfAddrspacePass());
if (LangOpts.EnableDAEInSpirKernels)
MPM.addPass(DeadArgumentEliminationSYCLPass());

// Add SPIRITTAnnotations pass to the pass manager if
// -fsycl-instrument-device-code option was passed. This option can be
// used only with spir triple.
if (CodeGenOpts.SPIRITTAnnotations) {
assert(
TargetTriple.isSPIR() &&
"ITT annotations can only be added to a module with spir target");
MPM.addPass(SPIRITTAnnotationsPass());
}

// Allocate static local memory in SYCL kernel scope for each allocation
// call. It should be called after inlining pass.
if (LangOpts.SYCLIsDevice) {
// Group local memory pass depends on inlining. Turn it on even in case if
// all llvm passes or SYCL early optimizations are disabled.
// FIXME: Remove this workaround when dependency on inlining is eliminated.
if (CodeGenOpts.DisableLLVMPasses)
MPM.addPass(AlwaysInlinerPass(false));
MPM.addPass(SYCLLowerWGLocalMemoryPass());
// Allocate static local memory in SYCL kernel scope for each allocation
// call.
MPM.addPass(SYCLLowerWGLocalMemoryPass());
}
}

// Add a verifier pass if requested. We don't have to do this if the action
Expand Down
5 changes: 0 additions & 5 deletions clang/lib/Frontend/CompilerInvocation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1692,11 +1692,6 @@ bool CompilerInvocation::ParseCodeGenArgs(CodeGenOptions &Opts, ArgList &Args,
{std::string(Split.first), std::string(Split.second)});
}

Opts.DisableLLVMPasses =
Args.hasArg(OPT_disable_llvm_passes) ||
(Args.hasArg(OPT_fsycl_is_device) && T.isSPIR() &&
Args.hasArg(OPT_fno_sycl_early_optimizations));

const llvm::Triple::ArchType DebugEntryValueArchs[] = {
llvm::Triple::x86, llvm::Triple::x86_64, llvm::Triple::aarch64,
llvm::Triple::arm, llvm::Triple::armeb, llvm::Triple::mips,
Expand Down
18 changes: 6 additions & 12 deletions clang/test/CodeGenSYCL/device_has.cpp
Original file line number Diff line number Diff line change
@@ -1,38 +1,32 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

// Tests for IR of device_has(aspect, ...) attribute and
// !sycl_used_aspects metadata
// Tests for IR of device_has(aspect, ...) attribute
#include "sycl.hpp"

using namespace sycl;
queue q;

// CHECK: define dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]]

// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]] {
[[sycl::device_has(sycl::aspect::cpu)]] void func1() {}

// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS2]]
// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]] {
[[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {}

// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
[[sycl::device_has()]] void func3() {}

// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS3]]
// CHECK: define linkonce_odr spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]] {
template <sycl::aspect Aspect>
[[sycl::device_has(Aspect)]] void func4() {}

// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]] {
[[sycl::device_has(sycl::aspect::cpu)]] void func5();
void func5() {}

constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]] {
[[sycl::device_has(getAspect())]] void func6() {}

class KernelFunctor {
Expand Down
11 changes: 8 additions & 3 deletions clang/test/CodeGenSYCL/group-local-memory.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -12,13 +12,18 @@
// Check that AlwaysInliner pass is always run for compilation of SYCL device
// target code, even if all optimizations are disabled.

// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm -disable-llvm-passes \
// RUN: -mdebug-pass Structure %s -o /dev/null 2>&1 \
// RUN: | FileCheck %s --check-prefixes=CHECK-ALWINL,CHECK
// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm -fno-sycl-early-optimizations \
// RUN: -mdebug-pass Structure %s -o /dev/null 2>&1 \
// RUN: | FileCheck %s --check-prefixes=CHECK-ALWINL,CHECK

// CHECK-INL: Running pass: ModuleInlinerWrapperPass on [module]
// CHECK-ALWINL: Running pass: AlwaysInlinerPass on [module]
// CHECK: Running pass: SYCLLowerWGLocalMemoryPass on [module]

// RUN: %clang_cc1 -fsycl-is-device -triple spir64-unknown-unknown -emit-llvm -disable-llvm-passes \
// RUN: -mdebug-pass Structure %s -o /dev/null 2>&1 \
// RUN: | FileCheck %s --check-prefixes=CHECK-NO-PASSES-ALWINL,CHECK-NO-PASSES,CHECK-NO-PASSES-INL

// CHECK-NO-PASSES-INL-NOT: Running pass: ModuleInlinerWrapperPass on [module]
// CHECK-NO-PASSES-ALWINL-NOT: Running pass: AlwaysInlinerPass on [module]
// CHECK-NO-PASSES-NOT: Running pass: SYCLLowerWGLocalMemoryPass on [module]
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/sub-group-size.cpp
Original file line number Diff line number Diff line change
@@ -1,6 +1,6 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=NONE,ALL
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=primary -sycl-std=2020 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=PRIM_DEF,ALL
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=10 -sycl-std=2020 -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s --check-prefixes=TEN_DEF,ALL
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64-unknown-unknown -fno-sycl-early-optimizations -emit-llvm -o - %s | FileCheck %s --check-prefixes=NONE,ALL
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=primary -sycl-std=2020 -triple spir64-unknown-unknown -fno-sycl-early-optimizations -emit-llvm -o - %s | FileCheck %s --check-prefixes=PRIM_DEF,ALL
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsycl-default-sub-group-size=10 -sycl-std=2020 -triple spir64-unknown-unknown -fno-sycl-early-optimizations -emit-llvm -o - %s | FileCheck %s --check-prefixes=TEN_DEF,ALL

// Ensure that both forms of the new sub_group_size properly emit their metadata
// on sycl-kernel and sycl-external functions.
Expand Down
15 changes: 7 additions & 8 deletions clang/test/CodeGenSYCL/uses_aspects.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -28,7 +28,7 @@ template <sycl::aspect Aspect>
void func5() {}

[[__sycl_detail__::__uses_aspects__(sycl::aspect::cpu)]] void func6();
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_used_aspects ![[ASPECTS4:[0-9]+]] {
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_used_aspects ![[ASPECTS1]] {
void func6() {
Type1WithAspect TestObj1;
Type2WithAspect TestObj2;
Expand Down Expand Up @@ -58,10 +58,9 @@ void foo() {
});
}
// CHECK: !sycl_types_that_use_aspects = !{![[TYPE1:[0-9]+]], ![[TYPE2:[0-9]+]]}
// CHECK-DAG: [[TYPE1]] = !{!"class.Type1WithAspect", i32 1}
// CHECK-DAG: [[TYPE2]] = !{!"class.Type2WithAspect", i32 5, i32 1}
// CHECK-DAG: [[EMPTYASPECTS]] = !{}
// CHECK-DAG: [[ASPECTS1]] = !{i32 1}
// CHECK-DAG: [[ASPECTS2]] = !{i32 5, i32 2}
// CHECK-DAG: [[ASPECTS3]] = !{i32 0}
// CHECK-DAG: [[ASPECTS4]] = !{i32 1, i32 5}
// CHECK: [[TYPE1]] = !{!"class.Type1WithAspect", i32 1}
// CHECK: [[TYPE2]] = !{!"class.Type2WithAspect", i32 5, i32 1}
// CHECK: [[EMPTYASPECTS]] = !{}
// CHECK: [[ASPECTS1]] = !{i32 1}
// CHECK: [[ASPECTS2]] = !{i32 5, i32 2}
// CHECK: [[ASPECTS3]] = !{i32 0}
43 changes: 43 additions & 0 deletions clang/test/SemaSYCL/sycl-force-inline-kernel-lambda-ast.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,43 @@
// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -ast-dump -o - %s | FileCheck %s --check-prefixes=NOINLINE,CHECK
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown -ast-dump -o - %s | FileCheck %s --check-prefixes=INLINE,CHECK

// Tests that the appropriate inlining attributes are added to kernel lambda functions,
// with no inline attribute being added when -fno-sycl-force-inline-kernel-lambda is set
// and attribute not explicitly provided.

#include "sycl.hpp"

int main() {
sycl::queue q;

q.submit([&](sycl::handler &h) {
// CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:17
// INLINE: AlwaysInlineAttr
// NOINLINE-NOT: AlwaysInlineAttr
h.parallel_for<class KernelName>([] {});
});

q.submit([&](sycl::handler &h) {
// CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:23
// CHECK: AlwaysInlineAttr
h.parallel_for<class KernelNameInline>([]() __attribute__((always_inline)) {});
});

q.submit([&](sycl::handler &h) {
// CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:30
// CHECK: NoInlineAttr
// CHECK-NOT: AlwaysInlineAttr
h.parallel_for<class KernelNameNoInline>([]() __attribute__((noinline)) {});
});

/// The flag is ignored for ESIMD kernels
q.submit([&](sycl::handler &h) {
// CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:39
// CHECK: SYCLSimdAttr
// CHECK-NOT: AlwaysInlineAttr
// CHECK-NOT: NoInlineAttr
h.parallel_for<class KernelNameESIMD>([]() __attribute__((sycl_explicit_simd)) {});
});

return 0;
}
30 changes: 0 additions & 30 deletions clang/test/SemaSYCL/sycl-force-inline-kernel-lambda.cpp

This file was deleted.

71 changes: 71 additions & 0 deletions sycl/test/check_device_code/device_has.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,71 @@
// RUN: %clangxx -fsycl -Xclang -fsycl-is-device -fsycl-device-only -Xclang -fno-sycl-early-optimizations -S -emit-llvm %s -o - | FileCheck %s

// Tests for IR of device_has(aspect, ...) attribute and
// !sycl_used_aspects metadata
#include <sycl/sycl.hpp>

using namespace sycl;
queue q;

// CHECK: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_1{{.*}} !sycl_declared_aspects ![[ASPECTS1:[0-9]+]] !srcloc ![[SRCLOC1:[0-9]+]] {{.*}}

// CHECK: define dso_local spir_func void @{{.*}}func1{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC2:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
[[sycl::device_has(sycl::aspect::cpu)]] void func1() {}

// CHECK: define dso_local spir_func void @{{.*}}func2{{.*}} !sycl_declared_aspects ![[ASPECTS2:[0-9]+]] !srcloc ![[SRCLOC3:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS2]]
[[sycl::device_has(sycl::aspect::fp16, sycl::aspect::gpu)]] void func2() {}

// CHECK: define dso_local spir_func void @{{.*}}func3{{.*}} !sycl_declared_aspects ![[EMPTYASPECTS:[0-9]+]] !srcloc ![[SRCLOC4:[0-9]+]] {
[[sycl::device_has()]] void func3() {}

// CHECK: define linkonce_odr dso_local spir_func void @{{.*}}func4{{.*}} !sycl_declared_aspects ![[ASPECTS3:[0-9]+]] !srcloc ![[SRCLOC5:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS3]]
template <sycl::aspect Aspect> [[sycl::device_has(Aspect)]] void func4() {}

// CHECK: define dso_local spir_func void @{{.*}}func5{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC6:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
[[sycl::device_has(sycl::aspect::cpu)]] void func5();
void func5() {}

constexpr sycl::aspect getAspect() { return sycl::aspect::cpu; }
// CHECK: define dso_local spir_func void @{{.*}}func6{{.*}} !sycl_declared_aspects ![[ASPECTS1]] !srcloc ![[SRCLOC7:[0-9]+]]
// CHECK-SAME: !sycl_used_aspects ![[ASPECTS1]]
[[sycl::device_has(getAspect())]] void func6() {}

class KernelFunctor {
public:
[[sycl::device_has(sycl::aspect::cpu)]] void operator()() const {
func1();
func2();
func3();
func4<sycl::aspect::host>();
func5();
func6();
}
};

void foo() {
q.submit([&](handler &h) {
KernelFunctor f1;
h.single_task<class kernel_name_1>(f1);
// CHECK: define weak_odr dso_local spir_kernel void @{{.*}}kernel_name_2{{.*}} !sycl_declared_aspects ![[ASPECTS4:[0-9]+]] !srcloc ![[SRCLOC8:[0-9]+]] {{.*}}
h.single_task<class kernel_name_2>(
[]() [[sycl::device_has(sycl::aspect::gpu)]] {});
});
}

// CHECK: [[ASPECTS1]] = !{i32 1}
// CHECK: [[SRCLOC1]] = !{i32 {{[0-9]+}}}
// CHECK: [[EMPTYASPECTS]] = !{}
// CHECK: [[SRCLOC2]] = !{i32 {{[0-9]+}}}
// CHECK: [[ASPECTS2]] = !{i32 5, i32 2}
// CHECK: [[SRCLOC3]] = !{i32 {{[0-9]+}}}
// CHECK: [[SRCLOC4]] = !{i32 {{[0-9]+}}}
// CHECK: [[ASPECTS3]] = !{i32 0}
// CHECK: [[SRCLOC5]] = !{i32 {{[0-9]+}}}
// CHECK: [[SRCLOC6]] = !{i32 {{[0-9]+}}}
// CHECK: [[SRCLOC7]] = !{i32 {{[0-9]+}}}
// CHECK: [[ASPECTS4]] = !{i32 2}
// CHECK: [[SRCLOC8]] = !{i32 {{[0-9]+}}}