Skip to content

Commit 92c23b1

Browse files
steffenlarsenmaksimsab
steffenlarsen
and
maksimsab
authored
[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 <[email protected]> Co-authored-by: Sabianin, Maksim <[email protected]>
1 parent 5bb9ef4 commit 92c23b1

23 files changed

+1177
-111
lines changed

‎clang/include/clang/Basic/DiagnosticFrontendKinds.td‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -281,6 +281,13 @@ def warn_avx_calling_convention
281281
InGroup<DiagGroup<"psabi">>;
282282
def err_avx_calling_convention : Error<warn_avx_calling_convention.Text>;
283283

284+
def warn_sycl_device_has_aspect_mismatch
285+
: Warning<"function '%0' uses aspect '%1' not listed in its "
286+
"'sycl::device_has' attribute">, BackendInfo,
287+
InGroup<SyclAspectMismatch>;
288+
def note_sycl_aspect_propagated_from_call
289+
: Note<"propagated from call to function '%0'">, BackendInfo;
290+
284291
def err_alias_to_undefined : Error<
285292
"%select{alias|ifunc}0 must point to a defined "
286293
"%select{variable or |}1function">;

‎clang/include/clang/Basic/DiagnosticGroups.td‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1277,6 +1277,7 @@ def Sycl2020Compat : DiagGroup<"sycl-2020-compat">;
12771277
def SyclStrict : DiagGroup<"sycl-strict", [ Sycl2017Compat, Sycl2020Compat]>;
12781278
def SyclTarget : DiagGroup<"sycl-target">;
12791279
def SyclFPGAMismatch : DiagGroup<"sycl-fpga-mismatch">;
1280+
def SyclAspectMismatch : DiagGroup<"sycl-aspect-mismatch">;
12801281

12811282
// Backend warnings.
12821283
def BackendInlineAsm : DiagGroup<"inline-asm">;

‎clang/lib/CodeGen/BackendUtil.cpp‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,7 @@
4747
#include "llvm/SYCLLowerIR/ESIMD/ESIMDVerifier.h"
4848
#include "llvm/SYCLLowerIR/LowerWGLocalMemory.h"
4949
#include "llvm/SYCLLowerIR/MutatePrintfAddrspace.h"
50+
#include "llvm/SYCLLowerIR/SYCLPropagateAspectsUsage.h"
5051
#include "llvm/Support/BuryPointer.h"
5152
#include "llvm/Support/CommandLine.h"
5253
#include "llvm/Support/MemoryBuffer.h"
@@ -876,6 +877,11 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
876877

877878
ModulePassManager MPM;
878879

880+
// FIXME: Change this when -fno-sycl-early-optimizations is not tied to
881+
// -disable-llvm-passes.
882+
if (CodeGenOpts.DisableLLVMPasses && LangOpts.SYCLIsDevice)
883+
MPM.addPass(SYCLPropagateAspectsUsagePass());
884+
879885
if (!CodeGenOpts.DisableLLVMPasses) {
880886
// Map our optimization levels into one of the distinct levels used to
881887
// configure the pipeline.
@@ -885,6 +891,7 @@ void EmitAssemblyHelper::RunOptimizationPipeline(
885891
PB.registerPipelineStartEPCallback(
886892
[&](ModulePassManager &MPM, OptimizationLevel Level) {
887893
MPM.addPass(ESIMDVerifierPass(LangOpts.SYCLESIMDForceStatelessMem));
894+
MPM.addPass(SYCLPropagateAspectsUsagePass());
888895
});
889896

890897
bool IsThinLTO = CodeGenOpts.PrepareForThinLTO;

‎clang/lib/CodeGen/CodeGenAction.cpp‎

Lines changed: 22 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -463,6 +463,8 @@ namespace clang {
463463
/// Specialized handler for misexpect warnings.
464464
/// Note that misexpect remarks are emitted through ORE
465465
void MisExpectDiagHandler(const llvm::DiagnosticInfoMisExpect &D);
466+
void
467+
AspectMismatchDiagHandler(const llvm::DiagnosticInfoAspectsMismatch &D);
466468
};
467469

468470
void BackendConsumer::anchor() {}
@@ -858,6 +860,23 @@ void BackendConsumer::DontCallDiagHandler(const DiagnosticInfoDontCall &D) {
858860
<< llvm::demangle(D.getFunctionName().str()) << D.getNote();
859861
}
860862

863+
void BackendConsumer::AspectMismatchDiagHandler(
864+
const DiagnosticInfoAspectsMismatch &D) {
865+
SourceLocation LocCookie =
866+
SourceLocation::getFromRawEncoding(D.getLocCookie());
867+
assert(LocCookie.isValid() &&
868+
"Invalid location for caller in aspect mismatch diagnostic");
869+
Diags.Report(LocCookie, diag::warn_sycl_device_has_aspect_mismatch)
870+
<< llvm::demangle(D.getFunctionName().str()) << D.getAspect();
871+
for (const std::pair<StringRef, unsigned> &CalleeInfo : D.getCallChain()) {
872+
LocCookie = SourceLocation::getFromRawEncoding(CalleeInfo.second);
873+
assert(LocCookie.isValid() &&
874+
"Invalid location for callee in aspect mismatch diagnostic");
875+
Diags.Report(LocCookie, diag::note_sycl_aspect_propagated_from_call)
876+
<< llvm::demangle(CalleeInfo.first.str());
877+
}
878+
}
879+
861880
void BackendConsumer::MisExpectDiagHandler(
862881
const llvm::DiagnosticInfoMisExpect &D) {
863882
StringRef Filename;
@@ -959,6 +978,9 @@ void BackendConsumer::DiagnosticHandlerImpl(const DiagnosticInfo &DI) {
959978
case llvm::DK_MisExpect:
960979
MisExpectDiagHandler(cast<DiagnosticInfoMisExpect>(DI));
961980
return;
981+
case llvm::DK_AspectMismatch:
982+
AspectMismatchDiagHandler(cast<DiagnosticInfoAspectsMismatch>(DI));
983+
return;
962984
default:
963985
// Plugin IDs are not bound to any value as they are set dynamically.
964986
ComputeDiagRemarkID(Severity, backend_plugin, DiagID);

‎clang/lib/CodeGen/CodeGenFunction.cpp‎

Lines changed: 8 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1086,6 +1086,14 @@ void CodeGenFunction::StartFunction(GlobalDecl GD, QualType RetTy,
10861086
Fn->setMetadata("sycl_used_aspects",
10871087
llvm::MDNode::get(getLLVMContext(), AspectsMD));
10881088
}
1089+
1090+
// Source location of functions is required to emit required diagnostics in
1091+
// SYCLPropagateAspectsUsagePass. Save the token in a srcloc metadata node.
1092+
llvm::ConstantInt *Line =
1093+
llvm::ConstantInt::get(Int32Ty, D->getLocation().getRawEncoding());
1094+
llvm::ConstantAsMetadata *SrcLocMD = llvm::ConstantAsMetadata::get(Line);
1095+
llvm::MDTuple *SrcLocMDT = llvm::MDNode::get(getLLVMContext(), {SrcLocMD});
1096+
Fn->setMetadata("srcloc", SrcLocMDT);
10891097
}
10901098

10911099
if (getLangOpts().SYCLIsDevice && D &&

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -25,7 +25,7 @@ struct S {
2525
// CHECK-NEXT: br label [[COND_END]]
2626
// CHECK: cond.end:
2727
// CHECK-NEXT: [[COND_LVALUE:%.*]] = phi ptr addrspace(4) [ [[TMP1]], [[COND_TRUE]] ], [ [[RHS_ASCAST]], [[COND_FALSE]] ]
28-
// 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
28+
// 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]+}}
2929
// CHECK-NEXT: ret void
3030
//
3131
S foo(bool cond, S &lhs, S rhs) {

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

Lines changed: 34 additions & 34 deletions
Original file line numberDiff line numberDiff line change
@@ -159,177 +159,177 @@ class Functor11 {
159159

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

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

0 commit comments

Comments
 (0)