Skip to content

[SYCL] Rename intel_reqd_sub_group_size attribute metadata #7713

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

Closed
wants to merge 4 commits into from
Closed

[SYCL] Rename intel_reqd_sub_group_size attribute metadata #7713

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
f0ac8b9
[SYCL] Rename intel_reqd_sub_group_size attribute metadata
smanna12 Dec 8, 2022
3d69beb
update patch
smanna12 Dec 8, 2022
72d4b24
update patch
smanna12 Dec 8, 2022
cbeb554
update patch
smanna12 Dec 9, 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
12 changes: 6 additions & 6 deletions clang/lib/CodeGen/CodeGenFunction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -643,14 +643,14 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
Optional<llvm::APSInt> ArgVal = CE->getResultAsAPSInt();
llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get(
Builder.getInt32(ArgVal->getSExtValue()))};
Fn->setMetadata("intel_reqd_sub_group_size",
Fn->setMetadata("sycl_reqd_sub_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
} else if (IsKernelOrDevice &&
CGM.getLangOpts().getDefaultSubGroupSizeType() ==
LangOptions::SubGroupSizeType::Integer) {
llvm::Metadata *AttrMDArgs[] = {llvm::ConstantAsMetadata::get(
Builder.getInt32(CGM.getLangOpts().DefaultSubGroupSize))};
Fn->setMetadata("intel_reqd_sub_group_size",
Fn->setMetadata("sycl_reqd_sub_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
}

Expand All @@ -662,18 +662,18 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
Context, A->getType() == IntelNamedSubGroupSizeAttr::Primary
? "primary"
: "automatic")};
Fn->setMetadata("intel_reqd_sub_group_size",
Fn->setMetadata("sycl_reqd_sub_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
} else if (CGM.getLangOpts().getDefaultSubGroupSizeType() ==
LangOptions::SubGroupSizeType::Auto) {
llvm::Metadata *AttrMDArgs[] = {
llvm::MDString::get(Context, "automatic")};
Fn->setMetadata("intel_reqd_sub_group_size",
Fn->setMetadata("sycl_reqd_sub_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
} else if (CGM.getLangOpts().getDefaultSubGroupSizeType() ==
LangOptions::SubGroupSizeType::Primary) {
llvm::Metadata *AttrMDArgs[] = {llvm::MDString::get(Context, "primary")};
Fn->setMetadata("intel_reqd_sub_group_size",
Fn->setMetadata("sycl_reqd_sub_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
}
}
Expand All @@ -682,7 +682,7 @@ void CodeGenFunction::EmitKernelMetadata(const FunctionDecl *FD,
Fn->setMetadata("sycl_explicit_simd", llvm::MDNode::get(Context, {}));
llvm::Metadata *AttrMDArgs[] = {
llvm::ConstantAsMetadata::get(Builder.getInt32(1))};
Fn->setMetadata("intel_reqd_sub_group_size",
Fn->setMetadata("sycl_reqd_sub_group_size",
llvm::MDNode::get(Context, AttrMDArgs));
}

Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenOpenCL/kernel-attributes.cl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ kernel __attribute__((vec_type_hint(uint4))) __attribute__((work_group_size_hint
// CHECK: define {{(dso_local )?}}spir_kernel void @kernel2(i32 {{[^%]*}}%a) {{[^{]+}} !vec_type_hint ![[MD3:[0-9]+]] !work_group_size_hint ![[MD4:[0-9]+]]

kernel __attribute__((intel_reqd_sub_group_size(8))) void kernel3(int a) {}
// CHECK: define {{(dso_local )?}}spir_kernel void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !intel_reqd_sub_group_size ![[MD5:[0-9]+]]
// CHECK: define {{(dso_local )?}}spir_kernel void @kernel3(i32 {{[^%]*}}%a) {{[^{]+}} !sycl_reqd_sub_group_size ![[MD5:[0-9]+]]
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm concerned about this change. This is an existing OpenCL C codegen test, why do we change the metadata name there? OpenCL flow is not being customized in our SYCL/DPC++ compiler as I understand it, so such change should go to upstream first, I suppose.

In OpenCL world, required sub-group size is an Intel extension and the translator is currently designed to only recognize intel_* metadata. By changing its name we will effectively break support for the attribute. See corresponding test and related commit in the translator repo.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I believe this change is a result of this discussion - #6804 (comment)

It isn't clear to me if the consensus was to actually modify this metadata. I don't see why we need to since this is an accepted Intel extension for openCL, and it exists outside of DPCPP. I understand the need to be consistent across various SYCL IR metadata but if SYCL is reusing existing constructs, I don't think we need to rename them. If we choose to do so, I agree with @AlexeySachkov that it should be done upstream.

I also can't recall why we have sycl::reqd_sub_group_size , intel::reqd_sub_group_size, intel_reqd_sub_group_size and intel::sub_group_size versions of this attribute. Can someone please explain?


// CHECK: [[MD1]] = !{i32 undef, i32 1}
// CHECK: [[MD2]] = !{i32 1, i32 2, i32 4}
Expand Down
10 changes: 5 additions & 5 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 @@ -239,15 +239,15 @@ int main() {
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]]{{.*}} !sycl_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]]{{.*}} !sycl_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]]{{.*}} !sycl_reqd_sub_group_size ![[NUM2]]
Functor5<2> f5;
h.single_task<class kernel_name19>(f5);

Expand Down Expand Up @@ -307,11 +307,11 @@ int main() {
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{{.*}} !sycl_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{{.*}} !sycl_reqd_sub_group_size ![[NUM1]]{{.*}} !sycl_explicit_simd ![[NUM]]
h.single_task<class kernel_name31>(
[]() [[intel::sycl_explicit_simd]]{});

Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/esimd_metadata1.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@

// The test checks that:
// 1. !sycl_explicit_simd metadata is generated for functions
// 2. !intel_reqd_sub_group_size !1 is added to explicit SIMD
// 2. !sycl_reqd_sub_group_size !1 is added to explicit SIMD
// kernel
// 3. Proper module !spirv.Source metadata is generated

Expand All @@ -15,10 +15,10 @@ void kernel(const Func &f) __attribute__((sycl_kernel)) {

void bar() {
kernel<class MyKernel>([=]() __attribute__((sycl_explicit_simd)){});
// CHECK: define {{.*}}spir_kernel void @_ZTSZ3barvE8MyKernel() {{.*}} !sycl_explicit_simd ![[EMPTY:[0-9]+]] !intel_reqd_sub_group_size ![[REQD_SIZE:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @_ZTSZ3barvE8MyKernel() {{.*}} !sycl_explicit_simd ![[EMPTY:[0-9]+]] !sycl_reqd_sub_group_size ![[REQD_SIZE:[0-9]+]]

kernel<class MyEsimdKernel>([=]() [[intel::sycl_explicit_simd]]{});
// CHECK: define {{.*}}spir_kernel void @_ZTSZ3barvE13MyEsimdKernel() {{.*}} !sycl_explicit_simd ![[EMPTY:[0-9]+]] !intel_reqd_sub_group_size ![[REQD_SIZE]]
// CHECK: define {{.*}}spir_kernel void @_ZTSZ3barvE13MyEsimdKernel() {{.*}} !sycl_explicit_simd ![[EMPTY:[0-9]+]] !sycl_reqd_sub_group_size ![[REQD_SIZE]]
}

// CHECK: !spirv.Source = !{[[LANG:![0-9]+]]}
Expand Down
6 changes: 3 additions & 3 deletions clang/test/CodeGenSYCL/esimd_metadata2.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,14 +1,14 @@
// RUN: %clang_cc1 -disable-llvm-passes -triple spir64-unknown-unknown -fsycl-is-device -S -emit-llvm %s -o - | FileCheck %s --check-prefixes CHECK,CHECK-ESIMD

// This test checks that attribute !intel_reqd_sub_group_size !1
// This test checks that attribute !sycl_reqd_sub_group_size !1
// is added for kernels with !sycl_explicit_simd

void shared_func() { }

__attribute__((sycl_device)) __attribute__((sycl_explicit_simd)) void esimd_func() { shared_func(); }

// CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[SGSIZE1]] {{.*}}{
// CHECK-ESIMD-DAG: define {{.*}}spir_kernel void @{{.*}}kernel_cm() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[SGSIZE1:[0-9]+]] {{.*}}{
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}esimd_funcv() #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[SGSIZE1]] {{.*}}{
// CHECK-ESIMD-DAG: define {{.*}}spir_func void @{{.*}}shared_funcv() #{{[0-9]+}}{{.*}} {
// CHECK-ESIMD-DAG: define linkonce_odr spir_func void @_ZN12ESIMDFunctorclEv({{.*}}) #{{[0-9]+}}{{.*}} !sycl_explicit_simd !{{[0-9]+}} {{.*}}{

Expand Down
10 changes: 5 additions & 5 deletions clang/test/CodeGenSYCL/no_opaque_check-direct-attribute-propagation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -239,15 +239,15 @@ int main() {
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]]{{.*}} !sycl_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]]{{.*}} !sycl_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]]{{.*}} !sycl_reqd_sub_group_size ![[NUM2]]
Functor5<2> f5;
h.single_task<class kernel_name19>(f5);

Expand Down Expand Up @@ -307,11 +307,11 @@ int main() {
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{{.*}} !sycl_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{{.*}} !sycl_reqd_sub_group_size ![[NUM1]]{{.*}} !sycl_explicit_simd ![[NUM]]
h.single_task<class kernel_name31>(
[]() [[intel::sycl_explicit_simd]]{});

Expand Down
2 changes: 1 addition & 1 deletion clang/test/CodeGenSYCL/reqd-sub-group-size-spirv-intrin.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@ int main() {
return 0;
}

// CHECK: define dso_local spir_kernel void @{{.*}}main{{.*}}kernel_class() {{.*}} !intel_reqd_sub_group_size ![[SUBGROUPSIZE:[0-9]+]]
// CHECK: define dso_local spir_kernel void @{{.*}}main{{.*}}kernel_class() {{.*}} !sycl_reqd_sub_group_size ![[SUBGROUPSIZE:[0-9]+]]
// CHECK: tail call spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}})

// CHECK: declare spir_func void @{{.*}}__spirv_ControlBarrier{{.*}}({{.*}})
Expand Down
10 changes: 5 additions & 5 deletions clang/test/CodeGenSYCL/reqd-sub-group-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,11 +49,11 @@ int main() {
return 0;
}

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE16:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE8:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE4:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !sycl_reqd_sub_group_size ![[SGSIZE16:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !sycl_reqd_sub_group_size ![[SGSIZE8:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !sycl_reqd_sub_group_size ![[SGSIZE4:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !sycl_reqd_sub_group_size ![[SGSIZE2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !sycl_reqd_sub_group_size ![[SGSIZE2]]
// CHECK: ![[SGSIZE16]] = !{i32 16}
// CHECK: ![[SGSIZE8]] = !{i32 8}
// CHECK: ![[SGSIZE4]] = !{i32 4}
Expand Down
16 changes: 8 additions & 8 deletions clang/test/CodeGenSYCL/sub-group-size.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,33 +9,33 @@
using namespace sycl;

[[intel::named_sub_group_size(primary)]] SYCL_EXTERNAL void external_primary() {}
// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_primary{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[PRIMARY:[0-9]+]]
// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_primary{{.*}}() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[PRIMARY:[0-9]+]]

[[intel::sub_group_size(10)]] SYCL_EXTERNAL void external_10() {}
// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_10{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN:[0-9]+]]
// ALL-DAG: define {{.*}}spir_func void @{{.*}}external_10{{.*}}() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[TEN:[0-9]+]]

SYCL_EXTERNAL void external_default_behavior() {}
// NONE-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}} !srcloc !{{[0-9]+}} !sycl_fixed_targets !{{[0-9]+}} {
// PRIM_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[PRIMARY]]
// TEN_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN]]
// PRIM_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[PRIMARY]]
// TEN_DEF-DAG: define {{.*}}spir_func void @{{.*}}external_default_behavior{{.*}}() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[TEN]]

void default_behavior() {
kernel_single_task<class Kernel1>([]() {
});
}
// NONE-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}} !srcloc !{{[0-9]+}} !kernel_arg_buffer_location !{{[0-9]+}} !sycl_fixed_targets !{{[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]]
// PRIM_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[PRIMARY]]
// TEN_DEF-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel1() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[TEN]]

void primary() {
kernel_single_task<class Kernel2>([]() [[intel::named_sub_group_size(primary)]]{});
}
// ALL-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel2() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[PRIMARY]]
// ALL-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel2() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[PRIMARY]]

void ten() {
kernel_single_task<class Kernel3>([]() [[intel::sub_group_size(10)]]{});
}
// ALL-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel3() #{{[0-9]+}}{{.*}} !intel_reqd_sub_group_size ![[TEN]]
// ALL-DAG: define {{.*}}spir_kernel void @{{.*}}Kernel3() #{{[0-9]+}}{{.*}} !sycl_reqd_sub_group_size ![[TEN]]

// PRIM_DEF: ![[PRIMARY]] = !{!"primary"}
// TEN_DEF: ![[TEN]] = !{i32 10}
4 changes: 2 additions & 2 deletions clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,8 +42,8 @@ int main() {
return 0;
}

// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !reqd_work_group_size ![[WGSIZE1:[0-9]+]] !intel_reqd_sub_group_size ![[SGSIZE1:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !reqd_work_group_size ![[WGSIZE:[0-9]+]] !sycl_reqd_sub_group_size ![[SGSIZE:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !reqd_work_group_size ![[WGSIZE1:[0-9]+]] !sycl_reqd_sub_group_size ![[SGSIZE1:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]]
// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 {{.*}} !reqd_work_group_size ![[WGSIZE3:[0-9]+]]
// CHECK: ![[WGSIZE]] = !{i32 16, i32 16, i32 32}
Expand Down
2 changes: 1 addition & 1 deletion llvm/lib/SYCLLowerIR/LowerInvokeSimd.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,7 +39,7 @@ using namespace llvm::esimd;

namespace {

constexpr char REQD_SUB_GROUP_SIZE_MD[] = "intel_reqd_sub_group_size";
constexpr char REQD_SUB_GROUP_SIZE_MD[] = "sycl_reqd_sub_group_size";

class SYCLLowerInvokeSimdLegacyPass : public ModulePass {
public:
Expand Down
8 changes: 4 additions & 4 deletions llvm/test/tools/sycl-post-link/sycl-esimd/no-sycl-esimd-split-shared-func.ll
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,15 +40,15 @@ declare dso_local x86_regcallcc noundef float @_Z33__regcall3____builtin_invoke_

;---- This has linkonce_odr, so it should be removed after inlining into the
;---- helper.
define linkonce_odr x86_regcallcc <4 x float> @SIMD_CALLEE(<4 x float> %val) #3 !sycl_explicit_simd !0 !intel_reqd_sub_group_size !1 {
define linkonce_odr x86_regcallcc <4 x float> @SIMD_CALLEE(<4 x float> %val) #3 !sycl_explicit_simd !0 !sycl_reqd_sub_group_size !1 {
; CHECK-NOT: {{.*}} @SIMD_CALLEE(
%data = call spir_func <4 x float> @SHARED_F(i64 100)
%add = fadd <4 x float> %val, %data
ret <4 x float> %add
}

;---- Function containing the invoke_simd call.
define dso_local spir_func float @SPMD_CALLER(float %x) #0 !intel_reqd_sub_group_size !2 {
define dso_local spir_func float @SPMD_CALLER(float %x) #0 !sycl_reqd_sub_group_size !2 {
%res = call spir_func float @_Z33__regcall3____builtin_invoke_simdXX(<4 x float> (<4 x float> (<4 x float>)*, <4 x float>)* @SIMD_CALL_HELPER, <4 x float> (<4 x float>)* @SIMD_CALLEE, float %x)
ret float %res
}
Expand Down Expand Up @@ -114,9 +114,9 @@ attributes #3 = { "referenced-indirectly" }
;---- * %f removed
;---- * indirect call replaced with direct and inlined
;---- * linkonce_odr linkage replaced with weak_odr
;---- * sycl_explicit_simd and intel_reqd_sub_group_size attributes added, which
;---- * sycl_explicit_simd and sycl_reqd_sub_group_size attributes added, which
;---- is required for correct processing by LowerESIMD
; CHECK: define weak_odr dso_local x86_regcallcc <4 x float> @[[NEW_HELPER_NAME]](<4 x float> %{{.*}}) #[[NEW_HELPER_ATTRS:[0-9]+]] !sycl_explicit_simd !1 !intel_reqd_sub_group_size !2 {
; CHECK: define weak_odr dso_local x86_regcallcc <4 x float> @[[NEW_HELPER_NAME]](<4 x float> %{{.*}}) #[[NEW_HELPER_ATTRS:[0-9]+]] !sycl_explicit_simd !1 !sycl_reqd_sub_group_size !2 {
; CHECK-NEXT: %{{.*}} = call spir_func <4 x float> @SHARED_F.esimd(i64 100)
; CHECK-NEXT: %{{.*}} = fadd <4 x float> %{{.*}}, %{{.*}}
; CHECK-NEXT: ret <4 x float> {{.*}}
Expand Down
4 changes: 2 additions & 2 deletions sycl/doc/design/OptionalDeviceFeatures.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -581,7 +581,7 @@ each target device. The implementation looks for functions that have either
from these metadata to the allowed list in the configuration file. If any
aspect is not on the allowed list, the implementation issues a warning. In
addition, the implementation looks for device functions that have
`!intel_reqd_sub_group_size` or `!reqd_work_group_size`. If the required
`!sycl_reqd_sub_group_size` or `!reqd_work_group_size`. If the required
sub-group size or the required work-group size is not allowed for the target
devices according to the configuration file, the implementation issues a
warning.
Expand Down Expand Up @@ -771,7 +771,7 @@ If the image contains kernels that were *not* compiled with
`!sycl_declared_aspects` metadata and one of the aspects in that metadata
is not in the image's *UsedAspects* set, issue a warning and add that
aspect to the *FinalUsedAspects* set.
- If the function has `!intel_reqd_sub_group_size` metadata and the size is
- If the function has `!sycl_reqd_sub_group_size` metadata and the size is
not the same as the image's required sub-group size, issue a warning and
add that sub-group size to the *FinalSubGroup* set.
- (Since the `[[sycl::reqd_work_group_size()]]` attribute cannot be specified
Expand Down