diff --git a/clang/include/clang/Basic/CodeGenOptions.def b/clang/include/clang/Basic/CodeGenOptions.def index 10fc983578e1e..2042ee23702ec 100644 --- a/clang/include/clang/Basic/CodeGenOptions.def +++ b/clang/include/clang/Basic/CodeGenOptions.def @@ -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 diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index bffb025eedbc2..cf607985d57c4 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -5110,7 +5110,8 @@ def : Flag<["-"], "fno-sycl-explicit-simd">, Flags<[CoreOption, Deprecated]>, 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>; def fsycl_dead_args_optimization : Flag<["-"], "fsycl-dead-args-optimization">, Group, Flags<[NoArgumentUnused, CoreOption]>, HelpText<"Enables " "elimination of DPC++ dead kernel arguments">; diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp index 7f98ab65b43f2..05e66f57a3542 100644 --- a/clang/lib/CodeGen/BackendUtil.cpp +++ b/clang/lib/CodeGen/BackendUtil.cpp @@ -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. @@ -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); @@ -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 diff --git a/clang/lib/Frontend/CompilerInvocation.cpp b/clang/lib/Frontend/CompilerInvocation.cpp index a458a83ae9956..3548313d4f6d8 100644 --- a/clang/lib/Frontend/CompilerInvocation.cpp +++ b/clang/lib/Frontend/CompilerInvocation.cpp @@ -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, diff --git a/clang/test/CodeGenSYCL/device_has.cpp b/clang/test/CodeGenSYCL/device_has.cpp index 3b626a3516c76..2f7ea8f28e5eb 100644 --- a/clang/test/CodeGenSYCL/device_has.cpp +++ b/clang/test/CodeGenSYCL/device_has.cpp @@ -1,7 +1,6 @@ // 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; @@ -9,30 +8,25 @@ 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::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 { diff --git a/clang/test/CodeGenSYCL/group-local-memory.cpp b/clang/test/CodeGenSYCL/group-local-memory.cpp index ec7c3cb22b510..40613d51cb0ee 100644 --- a/clang/test/CodeGenSYCL/group-local-memory.cpp +++ b/clang/test/CodeGenSYCL/group-local-memory.cpp @@ -12,9 +12,6 @@ // 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 @@ -22,3 +19,11 @@ // 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] diff --git a/clang/test/CodeGenSYCL/sub-group-size.cpp b/clang/test/CodeGenSYCL/sub-group-size.cpp index 693f1885860b7..528c89b918c88 100644 --- a/clang/test/CodeGenSYCL/sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/sub-group-size.cpp @@ -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. diff --git a/clang/test/CodeGenSYCL/uses_aspects.cpp b/clang/test/CodeGenSYCL/uses_aspects.cpp index 00ff464ff941d..eb3ec5d815d79 100644 --- a/clang/test/CodeGenSYCL/uses_aspects.cpp +++ b/clang/test/CodeGenSYCL/uses_aspects.cpp @@ -28,7 +28,7 @@ template 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; @@ -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} diff --git a/clang/test/SemaSYCL/sycl-force-inline-kernel-lambda-ast.cpp b/clang/test/SemaSYCL/sycl-force-inline-kernel-lambda-ast.cpp new file mode 100644 index 0000000000000..13e8cffd6b911 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-force-inline-kernel-lambda-ast.cpp @@ -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([] {}); + }); + + q.submit([&](sycl::handler &h) { + // CHECK: LambdaExpr{{.*}}sycl-force-inline-kernel-lambda-ast.cpp:23 + // CHECK: AlwaysInlineAttr + h.parallel_for([]() __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([]() __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([]() __attribute__((sycl_explicit_simd)) {}); + }); + + return 0; +} diff --git a/clang/test/SemaSYCL/sycl-force-inline-kernel-lambda.cpp b/clang/test/SemaSYCL/sycl-force-inline-kernel-lambda.cpp deleted file mode 100644 index c8be5ca3467b2..0000000000000 --- a/clang/test/SemaSYCL/sycl-force-inline-kernel-lambda.cpp +++ /dev/null @@ -1,30 +0,0 @@ -// RUN: %clang_cc1 -fno-sycl-force-inline-kernel-lambda -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unkown-unknown -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,CHECK-NO-INLINE -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unkown-unknown -emit-llvm -o - %s | FileCheck %s --check-prefixes=CHECK,CHECK-INLINE - -#include "sycl.hpp" - -int main() { - sycl::queue q; - - // CHECK: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_E10KernelName() - // - // CHECK-NO-INLINE: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv - // CHECK-INLINE-NOT: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_ENKUlvE_clEv - q.submit([&](sycl::handler &h) { h.parallel_for([] {}); }); - - - // CHECK: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_E16KernelNameInline() - // CHECK-NOT: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_ENKUlvE_clEv - q.submit([&](sycl::handler &h) { h.parallel_for([]() __attribute__((always_inline)) {}); }); - - // CHECK: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE1_clES2_E18KernelNameNoInline() - // CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE1_clES2_ENKUlvE_clEv - q.submit([&](sycl::handler &h) { h.parallel_for([]() __attribute__((noinline)) {}); }); - - /// The flag is ignored for ESIMD kernels - // CHECK: define {{.*}} spir_kernel void @_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE2_clES2_E15KernelNameESIMD() - // CHECK: call spir_func void @_ZZZ4mainENKUlRN4sycl3_V17handlerEE2_clES2_ENKUlvE_clEv - q.submit([&](sycl::handler &h) { h.parallel_for([]() __attribute__((sycl_explicit_simd)) {}); }); - - return 0; -} diff --git a/sycl/test/check_device_code/device_has.cpp b/sycl/test/check_device_code/device_has.cpp new file mode 100644 index 0000000000000..d570bc9d370a5 --- /dev/null +++ b/sycl/test/check_device_code/device_has.cpp @@ -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 + +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::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(); + func5(); + func6(); + } +}; + +void foo() { + q.submit([&](handler &h) { + KernelFunctor f1; + h.single_task(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( + []() [[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]+}}}