Skip to content

[SYCL] Enable aspect usage propagation pass and add diagnostics #6982

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Nov 4, 2022
Merged

[SYCL] Enable aspect usage propagation pass and add diagnostics #6982

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
9596d42
[SYCL] Implement optional device feature diagnostics
steffenlarsen Sep 22, 2022
821c661
Merge remote-tracking branch 'intel/sycl' into steffen/aspects_prop_d…
steffenlarsen Oct 17, 2022
1682594
Respond to comments
steffenlarsen Oct 17, 2022
c2b083f
Merge remote-tracking branch 'intel/sycl' into steffen/aspects_prop_d…
steffenlarsen Oct 27, 2022
3ff7463
Merge remote-tracking branch 'intel/sycl' into steffen/aspects_prop_d…
steffenlarsen Nov 2, 2022
d21892f
Fix sub-group-size after merge
steffenlarsen Nov 2, 2022
4869f03
Merge remote-tracking branch 'intel/sycl' into steffen/aspects_prop_d…
steffenlarsen Nov 3, 2022
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
7 changes: 7 additions & 0 deletions clang/include/clang/Basic/DiagnosticFrontendKinds.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -276,6 +276,13 @@ def warn_avx_calling_convention
InGroup<DiagGroup<"psabi">>;
def err_avx_calling_convention : Error<warn_avx_calling_convention.Text>;

def warn_sycl_device_has_aspect_mismatch
: Warning<"function '%0' uses aspect '%1' not listed in its "
"'sycl::device_has' attribute">, BackendInfo,
InGroup<SyclAspectMismatch>;
def note_sycl_aspect_propagated_from_call
: Note<"propagated from call to function '%0'">, BackendInfo;

def err_alias_to_undefined : Error<
"%select{alias|ifunc}0 must point to a defined "
"%select{variable or |}1function">;
Expand Down
1 change: 1 addition & 0 deletions clang/include/clang/Basic/DiagnosticGroups.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1275,6 +1275,7 @@ def Sycl2020Compat : DiagGroup<"sycl-2020-compat">;
def SyclStrict : DiagGroup<"sycl-strict", [ Sycl2017Compat, Sycl2020Compat]>;
def SyclTarget : DiagGroup<"sycl-target">;
def SyclFPGAMismatch : DiagGroup<"sycl-fpga-mismatch">;
def SyclAspectMismatch : DiagGroup<"sycl-aspect-mismatch">;

// Backend warnings.
def BackendInlineAsm : DiagGroup<"inline-asm">;
Expand Down
7 changes: 7 additions & 0 deletions clang/lib/CodeGen/BackendUtil.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -47,6 +47,7 @@
#include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
#include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
#include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
#include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
#include "llvm/Support/BuryPointer.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/MemoryBuffer.h"
Expand Down Expand Up @@ -875,6 +876,11 @@ 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 All @@ -884,6 +890,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
PB.registerPipelineStartEPCallback(
[&](ModulePassManager &MPM, OptimizationLevel Level) {
MPM.addPass(ESIMDVerifierPass(LangOpts.SYCLESIMDForceStatelessMem));
MPM.addPass(SYCLPropagateAspectsUsagePass());
});

bool IsThinLTO = CodeGenOpts.PrepareForThinLTO;
Expand Down
22 changes: 22 additions & 0 deletions clang/lib/CodeGen/CodeGenAction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -458,6 +458,8 @@ namespace clang {
/// Specialized handler for misexpect warnings.
/// Note that misexpect remarks are emitted through ORE
void MisExpectDiagHandler(const llvm::DiagnosticInfoMisExpect &D);
void
AspectMismatchDiagHandler(const llvm::DiagnosticInfoAspectsMismatch &D);
};

void BackendConsumer::anchor() {}
Expand Down Expand Up @@ -839,6 +841,23 @@ void BackendConsumer::DontCallDiagHandler(const DiagnosticInfoDontCall &D) {
<< llvm::demangle(D.getFunctionName().str()) << D.getNote();
}

void BackendConsumer::AspectMismatchDiagHandler(
const DiagnosticInfoAspectsMismatch &D) {
SourceLocation LocCookie =
SourceLocation::getFromRawEncoding(D.getLocCookie());
assert(LocCookie.isValid() &&
"Invalid location for caller in aspect mismatch diagnostic");
Diags.Report(LocCookie, diag::warn_sycl_device_has_aspect_mismatch)
<< llvm::demangle(D.getFunctionName().str()) << D.getAspect();
for (const std::pair<StringRef, unsigned> &CalleeInfo : D.getCallChain()) {
LocCookie = SourceLocation::getFromRawEncoding(CalleeInfo.second);
assert(LocCookie.isValid() &&
"Invalid location for callee in aspect mismatch diagnostic");
Diags.Report(LocCookie, diag::note_sycl_aspect_propagated_from_call)
<< llvm::demangle(CalleeInfo.first.str());
}
}

void BackendConsumer::MisExpectDiagHandler(
const llvm::DiagnosticInfoMisExpect &D) {
StringRef Filename;
Expand Down Expand Up @@ -935,6 +954,9 @@ void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) {
case llvm::DK_MisExpect:
MisExpectDiagHandler(cast<DiagnosticInfoMisExpect>(DI));
return;
case llvm::DK_AspectMismatch:
AspectMismatchDiagHandler(cast<DiagnosticInfoAspectsMismatch>(DI));
return;
default:
// Plugin IDs are not bound to any value as they are set dynamically.
ComputeDiagRemarkID(Severity, backend_plugin, DiagID);
Expand Down
8 changes: 8 additions & 0 deletions clang/lib/CodeGen/CodeGenFunction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1086,6 +1086,14 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
Fn->setMetadata("sycl_used_aspects",
llvm::MDNode::get(getLLVMContext(), AspectsMD));
}

// Source location of functions is required to emit required diagnostics in
// SYCLPropagateAspectsUsagePass. Save the token in a srcloc metadata node.
llvm::ConstantInt *Line =
llvm::ConstantInt::get(Int32Ty, D->getLocation().getRawEncoding());
llvm::ConstantAsMetadata *SrcLocMD = llvm::ConstantAsMetadata::get(Line);
llvm::MDTuple *SrcLocMDT = llvm::MDNode::get(getLLVMContext(), {SrcLocMD});
Fn->setMetadata("srcloc", SrcLocMDT);
}

if (getLangOpts().SYCLIsDevice && D &&
Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/address-space-cond-op.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ struct S {
// CHECK-NEXT: br label [[COND_END]]
// CHECK: cond.end:
// CHECK-NEXT: [[COND_LVALUE:%.*]] = phi ptr addrspace(4) [ [[TMP1]], [[COND_TRUE]] ], [ [[RHS_ASCAST]], [[COND_FALSE]] ]
// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 2 %agg.result, ptr addrspace(4) align 2 [[COND_LVALUE]], i64 2, i1 false), !tbaa.struct !9
// CHECK-NEXT: call void @llvm.memcpy.p4.p4.i64(ptr addrspace(4) align 2 %agg.result, ptr addrspace(4) align 2 [[COND_LVALUE]], i64 2, i1 false), !tbaa.struct !{{[0-9]+}}
// CHECK-NEXT: ret void
//
S foo(bool cond, S &lhs, S rhs) {
Expand Down
68 changes: 34 additions & 34 deletions clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -159,177 +159,177 @@ class Functor11 {

int main() {
q.submit([&](handler &h) {
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 !kernel_arg_buffer_location ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0{{.*}} !kernel_arg_buffer_location ![[NUM:[0-9]+]]{{.*}} !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]]
Foo boo;
h.single_task<class kernel_name1>(boo);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]]
h.single_task<class kernel_name2>(
[]() [[intel::scheduler_target_fmax_mhz(42)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
Functor<2> f;
h.single_task<class kernel_name3>(f);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !scheduler_target_fmax_mhz
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z3foov()
h.single_task<class kernel_name4>(
[]() { foo(); });

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM1]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !num_simd_work_items ![[NUM1]]
Foo1 boo1;
h.single_task<class kernel_name5>(boo1);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM42]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !num_simd_work_items ![[NUM42]]
h.single_task<class kernel_name6>(
[]() [[intel::num_simd_work_items(42)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM2]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !num_simd_work_items ![[NUM2]]
Functor1<2> f1;
h.single_task<class kernel_name7>(f1);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !num_simd_work_items
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo1v()
h.single_task<class kernel_name8>(
[]() { foo1(); });

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !no_global_work_offset ![[NUM:[0-9]+]]
Foo2 boo2;
h.single_task<class kernel_name9>(boo2);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name10() #0 {{.*}} ![[NUM0:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name10() #0{{.*}} {{.*}} ![[NUM0:[0-9]+]]
h.single_task<class kernel_name10>(
[]() [[intel::no_global_work_offset(0)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !no_global_work_offset ![[NUM]]
Functor2<1> f2;
h.single_task<class kernel_name11>(f2);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !no_global_work_offset
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo2v()
h.single_task<class kernel_name12>(
[]() { foo2(); });

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !max_global_work_dim ![[NUM1]]
Foo3 boo3;
h.single_task<class kernel_name13>(boo3);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !max_global_work_dim ![[NUM1]]
h.single_task<class kernel_name14>(
[]() [[intel::max_global_work_dim(1)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM2]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !max_global_work_dim ![[NUM2]]
Functor3<2> f3;
h.single_task<class kernel_name15>(f3);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !max_global_work_dim
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo3v()
h.single_task<class kernel_name16>(
[]() { foo3(); });

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM16:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !intel_reqd_sub_group_size ![[NUM16:[0-9]+]]
Foo4 boo4;
h.single_task<class kernel_name17>(boo4);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM1]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !intel_reqd_sub_group_size ![[NUM1]]
h.single_task<class kernel_name18>(
[]() [[sycl::reqd_sub_group_size(1)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM2]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !intel_reqd_sub_group_size ![[NUM2]]
Functor5<2> f5;
h.single_task<class kernel_name19>(f5);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !reqd_sub_group_size
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo4v()
Functor4 f4;
h.single_task<class kernel_name20>(f4);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM32:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !reqd_work_group_size ![[NUM32:[0-9]+]]
Foo5 boo5;
h.single_task<class kernel_name21>(boo5);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM88:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !reqd_work_group_size ![[NUM88:[0-9]+]]
h.single_task<class kernel_name22>(
[]() [[sycl::reqd_work_group_size(8, 8, 8)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM22:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !reqd_work_group_size ![[NUM22:[0-9]+]]
Functor7<2, 2, 2> f7;
h.single_task<class kernel_name23>(f7);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !reqd_work_group_size
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo5v()
Functor6 f6;
h.single_task<class kernel_name24>(f6);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM32]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !max_work_group_size ![[NUM32]]
Foo6 boo6;
h.single_task<class kernel_name25>(boo6);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM88]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !max_work_group_size ![[NUM88]]
h.single_task<class kernel_name26>(
[]() [[intel::max_work_group_size(8, 8, 8)]]{});

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM22]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]{{.*}} !max_work_group_size ![[NUM22]]
Functor9<2, 2, 2> f9;
h.single_task<class kernel_name27>(f9);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !max_work_group_size
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo6v()
Functor8 f8;
h.single_task<class kernel_name28>(f8);

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK-NOT: !sycl_explicit_simd
// CHECK-SAME: {
// CHECK: define {{.*}}spir_func void @{{.*}}foo7{{.*}} !sycl_explicit_simd ![[NUM]]
h.single_task<class kernel_name29>(
[]() { foo7(); });

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name30() #0 !intel_reqd_sub_group_size ![[NUM1]] !sycl_explicit_simd ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name30() #0{{.*}} !intel_reqd_sub_group_size ![[NUM1]]{{.*}} !sycl_explicit_simd ![[NUM]]
Foo7 boo7;
h.single_task<class kernel_name30>(boo7);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name31() #0 !intel_reqd_sub_group_size ![[NUM1]] !sycl_explicit_simd ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name31() #0{{.*}} !intel_reqd_sub_group_size ![[NUM1]]{{.*}} !sycl_explicit_simd ![[NUM]]
h.single_task<class kernel_name31>(
[]() [[intel::sycl_explicit_simd]]{});

// Test attribute is not propagated.
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(ptr addrspace(4) noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
// CHECK-NOT: noalias
// CHECK-SAME: {
// CHECK: define dso_local spir_func void @_Z4foo8v()
Functor10 f10;
h.single_task<class kernel_name32>(f10);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #3 comdat align 2
Foo8 boo8;
h.single_task<class kernel_name33>(boo8);

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0 !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0{{.*}} !kernel_arg_buffer_location ![[NUM]]
// CHECK: define {{.*}}spir_func void @{{.*}}(ptr addrspace(4) noalias noundef align 1 dereferenceable_or_null(1) %this) #4 align 2
h.single_task<class kernel_name34>(
[]() [[intel::kernel_args_restrict]]{});
Expand Down
Loading