Skip to content

[SYCL] Add new kernel-arg-runtime-aligned metadata. #5111

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.

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 21 commits into from
Jan 18, 2022
Merged
Show file tree
Hide file tree
Changes from 5 commits
Commits
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
23 changes: 21 additions & 2 deletions clang/lib/CodeGen/CodeGenModule.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1699,6 +1699,10 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
// MDNode for the intel_buffer_location attribute.
SmallVector<llvm::Metadata *, 8> argSYCLBufferLocationAttr;

// MDNode for listing SYCL kernel pointer arguments originating from
// accessors
SmallVector<llvm::Metadata *, 8> argSYCLKernelRuntimeAligned;

// MDNode for listing ESIMD kernel pointer arguments originating from
// accessors
SmallVector<llvm::Metadata *, 8> argESIMDAccPtrs;
Expand Down Expand Up @@ -1750,6 +1754,12 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(
ArgInfoAddressSpace(pointeeTy.getAddressSpace()))));

// Get address qualifier of SYCL kernel pointer parameter from
// accessors.
argSYCLKernelRuntimeAligned.push_back(
llvm::ConstantAsMetadata::get(CGF->Builder.getInt1(
ArgInfoAddressSpace(pointeeTy.getAddressSpace()))));

// Get argument type name.
std::string typeName = getTypeSpelling(pointeeTy) + "*";
std::string baseTypeName =
Expand All @@ -1775,6 +1785,9 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,
addressQuals.push_back(
llvm::ConstantAsMetadata::get(CGF->Builder.getInt32(AddrSpc)));

argSYCLKernelRuntimeAligned.push_back(
llvm::ConstantAsMetadata::get(CGF->Builder.getInt1(AddrSpc)));

// Get argument type name.
ty = isPipe ? ty->castAs<PipeType>()->getElementType() : ty;
std::string typeName = getTypeSpelling(ty);
Expand Down Expand Up @@ -1813,10 +1826,16 @@ void CodeGenModule::GenOpenCLArgMetadata(llvm::Function *Fn,

bool IsEsimdFunction = FD && FD->hasAttr<SYCLSimdAttr>();

if (LangOpts.SYCLIsDevice && !IsEsimdFunction)
if (LangOpts.SYCLIsDevice && !IsEsimdFunction) {
Fn->setMetadata("kernel_arg_buffer_location",
llvm::MDNode::get(VMContext, argSYCLBufferLocationAttr));
else {

// The value of any "kernel_arg_runtime_aligned" metadata element is 1 for
// any kernel arguments that corresponds to the base pointer of an accessor
// and 0 otherwise.
Fn->setMetadata("kernel_arg_runtime_aligned",
llvm::MDNode::get(VMContext, argSYCLKernelRuntimeAligned));
} else {
Fn->setMetadata("kernel_arg_addr_space",
llvm::MDNode::get(VMContext, addressQuals));
Fn->setMetadata("kernel_arg_access_qual",
Expand Down
62 changes: 31 additions & 31 deletions clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -146,15 +146,15 @@ class Functor10 {

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]+]] !kernel_arg_runtime_aligned ![[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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]]
Functor<2> f;
h.single_task<class kernel_name3>(f);

Expand All @@ -166,128 +166,128 @@ int main() {
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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]+]]
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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]] !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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// CHECK-NOT: !sycl_explicit_simd
// CHECK-SAME: {
// CHECK: define {{.*}}spir_func void @{{.*}}foo7{{.*}} !sycl_explicit_simd ![[NUM]]
Expand All @@ -303,20 +303,20 @@ int main() {
[]() [[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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class.Functor10 addrspace(4)* 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.Foo8 addrspace(4)* noalias 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]] !kernel_arg_runtime_aligned ![[NUM:[0-9]+]]
// CHECK: define {{.*}}spir_func void @{{.*}}(%class.anon{{.*}} addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #4 align 2
h.single_task<class kernel_name34>(
[]() [[intel::kernel_args_restrict]]{});
Expand Down
69 changes: 69 additions & 0 deletions clang/test/CodeGenSYCL/kernel-arg-accessor-pointer.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,69 @@
// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown -disable-llvm-passes -emit-llvm %s -o - | FileCheck %s

// This test checks if the metadata "kernel-arg-runtime-aligned"
// is generated if the kernel captures an accessor

#include "sycl.hpp"

using namespace cl::sycl;

queue q;

int main() {

using Accessor =
accessor<int, 1, access::mode::read_write, access::target::global_buffer>;
Accessor acc[2];

// kernel_A parameters : int*, sycl::range<1>, sycl::range<1>, sycl::id<1>,
// int*, sycl::range<1>, sycl::range<1>,sycl::id<1>
q.submit([&](cl::sycl::handler &h) {
h.single_task<class kernel_A>([=]() {
acc[1].use();
});
});

// kernel_B parameters : none
q.submit([&](cl::sycl::handler &h) {
h.single_task<class kernel_B>([=]() {
int result = 5;
});
});

int a = 10;

// kernel_C parameters : int
q.submit([&](cl::sycl::handler &h) {
h.single_task<class kernel_C>([=]() {
int x = a;
});
});
}

// Check kernel_A parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_A
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG1:%[a-zA-Z0-9_]+]],
// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE1:%[a-zA-Z0-9_]+_1]],
// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE1:%[a-zA-Z0-9_]+_2]],
// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET1:%[a-zA-Z0-9_]+_3]],
// CHECK-SAME: i32 addrspace(1)* [[MEM_ARG2:%[a-zA-Z0-9_]+_4]],
// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[ACC_RANGE2:%[a-zA-Z0-9_]+_6]],
// CHECK-SAME: %"struct.cl::sycl::range"* byval{{.*}}align 4 [[MEM_RANGE2:%[a-zA-Z0-9_]+_7]],
// CHECK-SAME: %"struct.cl::sycl::id"* byval{{.*}}align 4 [[OFFSET2:%[a-zA-Z0-9_]+_8]])
// CHECK-SAME: !kernel_arg_runtime_aligned !5

// Check kernel_B parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_B
// CHECK-SAME: !kernel_arg_runtime_aligned !13

// Check kernel_C parameters
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_C
// CHECK-SAME: i32 [[MEM_ARG1:%[a-zA-Z0-9_]+]]
// CHECK-SAME: !kernel_arg_runtime_aligned !15

// Check kernel-arg-runtime-aligned metadata.
// The value of any metadata element is 1 for any kernel arguments
// that corresponds to the base pointer of an accessor and 0 otherwise.
// CHECK: !5 = !{i1 true, i1 false, i1 false, i1 false, i1 true, i1 false, i1 false, i1 false}
// CHECK: !13 = !{}
// CHECK: !15 = !{i1 false}
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/sub-group-size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ void default_behavior() {
kernel_single_task<class Kernel1>([]() {
});
}
// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} {
// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} !kernel_arg_runtime_aligned !{{[0-9]+}} {
// PRIM_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !intel_reqd_sub_group_size ![[PRIMARY]]
// TEN_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !intel_reqd_sub_group_size ![[TEN]]

Expand Down