Skip to content

Commit a023fdb

Browse files
author
iclsrc
committed
Merge from 'sycl' to 'sycl-web'
2 parents 3d964ee + 9670beb commit a023fdb

File tree

186 files changed

+3094
-807
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

186 files changed

+3094
-807
lines changed

‎.github/workflows/sycl_post_commit.yml

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -13,6 +13,7 @@ on:
1313
- .github/workflows/sycl_linux_build_and_test.yml
1414
- .github/workflows/sycl_windows_build_and_test.yml
1515
- .github/workflows/sycl_macos_build_and_test.yml
16+
workflow_dispatch:
1617

1718
jobs:
1819
# This job generates matrix of tests for LLVM Test Suite

‎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/lib/Sema/SemaSYCL.cpp

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3497,10 +3497,10 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler {
34973497

34983498
bool enterArray(FieldDecl *FD, QualType ArrayType,
34993499
QualType ElementType) final {
3500-
uint64_t ArraySize = SemaRef.getASTContext()
3501-
.getAsConstantArrayType(ArrayType)
3502-
->getSize()
3503-
.getZExtValue();
3500+
const ConstantArrayType *CAT =
3501+
SemaRef.getASTContext().getAsConstantArrayType(ArrayType);
3502+
assert(CAT && "Should only be called on constant-size array.");
3503+
uint64_t ArraySize = CAT->getSize().getZExtValue();
35043504
addCollectionInitListExpr(ArrayType, ArraySize);
35053505
ArrayInfos.emplace_back(getFieldEntity(FD, ArrayType), 0);
35063506

‎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)