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

This commit adds a warning diagnostic for when there is a mismatch
between aspect usage propagated to a function and the function's
`device_has` attribute. Additionally, notes accompany the warning to
give the user a trace through to where the aspects originate. Since the
aspects are propagated in an LLVM pass, the warning is issued from
there, which means no source information is available by default. To
alleviate this, Clang CodeGen will now add srcloc metadata with an
encoded version of the source location to allow the backend to correctly
report the location for the various new diagnostics.

Additionally, this commit also adds the SYCLPropagateAspectsUsage pass
to the passes run for SYCL device code.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
Co-authored-by: Sabianin, Maksim <maksim.sabianin@intel.com>

Author: steffenlarsen

clang/include/clang/Basic/DiagnosticFrontendKinds.td

@@ -281,6 +281,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">;

clang/include/clang/Basic/DiagnosticGroups.td

@@ -1277,6 +1277,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">;

clang/lib/CodeGen/BackendUtil.cpp

@@ -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"
@@ -876,6 +877,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.
@@ -885,6 +891,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
       PB.registerPipelineStartEPCallback(
           [&](ModulePassManager &MPM, OptimizationLevel Level) {
             MPM.addPass(ESIMDVerifierPass(LangOpts.SYCLESIMDForceStatelessMem));
+            MPM.addPass(SYCLPropagateAspectsUsagePass());
           });
 
     bool IsThinLTO = CodeGenOpts.PrepareForThinLTO;

clang/lib/CodeGen/CodeGenAction.cpp

@@ -463,6 +463,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() {}
@@ -858,6 +860,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;
@@ -959,6 +978,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);

clang/lib/CodeGen/CodeGenFunction.cpp

@@ -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 &&

clang/test/CodeGenSYCL/address-space-cond-op.cpp

@@ -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) {

clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp

@@ -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]]{});