From ac489a598fcb53f6fa233f069cfacde3819034de Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Fri, 27 Oct 2023 10:04:12 +0100 Subject: [PATCH 1/5] [SYCL] [AMDGPU] Ignore incorrect sub-group size CDNA supports only 64 wave front size, for those GPUs set subgroup size to 64. Some GPUS support both 32 and 64, for those (and the rest) only allow 32. --- .../clang/Basic/DiagnosticSemaKinds.td | 4 ++++ clang/lib/Sema/SemaDeclAttr.cpp | 19 ++++++++++++++- .../SemaSYCL/reqd-sub-group-size-amd_32.cpp | 24 +++++++++++++++++++ .../SemaSYCL/reqd-sub-group-size-amd_64.cpp | 24 +++++++++++++++++++ 4 files changed, 70 insertions(+), 1 deletion(-) create mode 100644 clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp create mode 100644 clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index a4328d78c0ccf..4f8ac42992c5c 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3526,6 +3526,10 @@ def warn_dllimport_dropped_from_inline_function : Warning< def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed" " only on a function directly called from a SYCL kernel function; attribute ignored">, InGroup; +def warn_amd_reqd_sub_group_attribute_n + : Warning<"attribute argument %0 is invalid and will be ignored; AMD " + "requires sub_group size %1">, + InGroup; def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with" " exception specification; attribute ignored">, InGroup; diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 22fc155343b1d..96221c300f0ac 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4007,10 +4007,27 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, << CI << /*positive*/ 0; return; } - if (Context.getTargetInfo().getTriple().isNVPTX() && ArgVal != 32) { + auto &TI = Context.getTargetInfo(); + if (TI.getTriple().isNVPTX() && ArgVal != 32) { Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_cuda_n_32) << ArgVal.getSExtValue(); } + if (TI.getTriple().isAMDGPU()) { + const auto HasWaveFrontSize64 = + TI.getTargetOpts().FeatureMap["wavefrontsize64"]; + const auto HasWaveFrontSize32 = + TI.getTargetOpts().FeatureMap["wavefrontsize32"]; + + // CDNA supports only 64 wave front size, for those GPUs allow subgroup + // size of 64. Some GPUs support both 32 and 64, for those (and the rest) + // only allow 32. + const auto SupportedWaveFrontSize = + HasWaveFrontSize64 && !HasWaveFrontSize32 ? 64 : 32; + if (ArgVal != SupportedWaveFrontSize) { + Diag(E->getExprLoc(), diag::warn_amd_reqd_sub_group_attribute_n) + << ArgVal.getSExtValue() << SupportedWaveFrontSize; + } + } // Check to see if there's a duplicate attribute with different values // already applied to the declaration. diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp new file mode 100644 index 0000000000000..0fa615e474d3b --- /dev/null +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify %s + +// Sub-group size is optimized for 32, warn (and ignore the attribute) if the +// size exceeds 32. +#include "sycl.hpp" + +int main() { + + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); // expected-warning {{attribute argument 64 is invalid and will be ignored; AMD requires sub_group size 32}} + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; AMD requires sub_group size 32}} + }); + + return 0; +} diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp new file mode 100644 index 0000000000000..88dbb8a295a4d --- /dev/null +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify %s + +// Sub-group size is optimized for 32, warn (and ignore the attribute) if the +// size exceeds 32. +#include "sycl.hpp" + +int main() { + + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; AMD requires sub_group size 64}} + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; AMD requires sub_group size 64}} + }); + + return 0; +} From 7076bdd4316cb15aefab99d75926f0deb40fd671 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Wed, 1 Nov 2023 13:52:19 +0000 Subject: [PATCH 2/5] PR fixes --- clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp | 2 +- clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp index 0fa615e474d3b..7b133f060680b 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -internal-isystem %S/Inputs -std=c++2b -verify %s // Sub-group size is optimized for 32, warn (and ignore the attribute) if the -// size exceeds 32. +// size is not 32. #include "sycl.hpp" int main() { diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp index 88dbb8a295a4d..98d43526b05e3 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl-is-device -triple amdgcn-amd-amdhsa -target-cpu gfx90a -internal-isystem %S/Inputs -std=c++2b -verify %s -// Sub-group size is optimized for 32, warn (and ignore the attribute) if the -// size exceeds 32. +// Sub-group size is optimized for 64, warn (and ignore the attribute) if the +// size is not 64. #include "sycl.hpp" int main() { From e9e31ee5d767e5d05ed2b6a6e66528a22255b3e0 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Tue, 7 Nov 2023 04:41:21 -0500 Subject: [PATCH 3/5] Early return and merge warning --- clang/include/clang/Basic/DiagnosticSemaKinds.td | 10 +++------- clang/lib/Sema/SemaDeclAttr.cpp | 11 +++++++---- clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp | 4 ++-- clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp | 4 ++-- clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp | 2 +- 5 files changed, 15 insertions(+), 16 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index 4ad747c4ff9dc..c6ee48befa5ba 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -3439,10 +3439,6 @@ def err_attribute_argument_is_zero : Error< def warn_attribute_argument_n_negative : Warning< "%0 attribute parameter %1 is negative and will be ignored">, InGroup; -def warn_reqd_sub_group_attribute_cuda_n_32 - : Warning<"attribute argument %0 is invalid and will be ignored; CUDA " - "requires sub_group size 32">, - InGroup; def err_property_function_in_objc_container : Error< "use of Objective-C property in function nested in Objective-C " "container not supported, move function outside its container">; @@ -3541,9 +3537,9 @@ def warn_dllimport_dropped_from_inline_function : Warning< def warn_attribute_on_direct_kernel_callee_only : Warning<"%0 attribute allowed" " only on a function directly called from a SYCL kernel function; attribute ignored">, InGroup; -def warn_amd_reqd_sub_group_attribute_n - : Warning<"attribute argument %0 is invalid and will be ignored; AMD " - "requires sub_group size %1">, +def warn_reqd_sub_group_attribute_n + : Warning<"attribute argument %0 is invalid and will be ignored; %1 " + "requires sub_group size %2">, InGroup; def warn_nothrow_attribute_ignored : Warning<"'nothrow' attribute conflicts with" " exception specification; attribute ignored">, diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 60f4d27cc8059..6ddfcea78eb9f 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4009,8 +4009,9 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, } auto &TI = Context.getTargetInfo(); if (TI.getTriple().isNVPTX() && ArgVal != 32) { - Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_cuda_n_32) - << ArgVal.getSExtValue(); + Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n) + << ArgVal.getSExtValue() << TI.getTriple().getArchName() << 32; + return; } if (TI.getTriple().isAMDGPU()) { const auto HasWaveFrontSize64 = @@ -4024,8 +4025,10 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, const auto SupportedWaveFrontSize = HasWaveFrontSize64 && !HasWaveFrontSize32 ? 64 : 32; if (ArgVal != SupportedWaveFrontSize) { - Diag(E->getExprLoc(), diag::warn_amd_reqd_sub_group_attribute_n) - << ArgVal.getSExtValue() << SupportedWaveFrontSize; + Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n) + << ArgVal.getSExtValue() << TI.getTriple().getArchName() + << SupportedWaveFrontSize; + return; } } diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp index 7b133f060680b..0b84abfa6b059 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp @@ -9,7 +9,7 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); // expected-warning {{attribute argument 64 is invalid and will be ignored; AMD requires sub_group size 32}} + h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); // expected-warning {{attribute argument 64 is invalid and will be ignored; amdgcn requires sub_group size 32}} }); Q.submit([&](sycl::handler &h) { @@ -17,7 +17,7 @@ int main() { }); Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; AMD requires sub_group size 32}} + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 32}} }); return 0; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp index 98d43526b05e3..6d457a8e0809d 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp @@ -13,11 +13,11 @@ int main() { }); Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; AMD requires sub_group size 64}} + h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; amdgcn requires sub_group size 64}} }); Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; AMD requires sub_group size 64}} + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 64}} }); return 0; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp index 2d9087ddee31d..c7800186a6786 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-cuda.cpp @@ -9,7 +9,7 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; CUDA requires sub_group size 32}} + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; nvptx requires sub_group size 32}} }); Q.submit([&](sycl::handler &h) { From 0a72b7b8aa1779171f615b3df22b13c3566fb20d Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Tue, 7 Nov 2023 09:20:59 -0500 Subject: [PATCH 4/5] Add CodeGen test to make sure that attributes are indeed ignored --- .../reqd-sub-group-size_ignored_values.cpp | 52 +++++++++++++++++++ .../SemaSYCL/reqd-sub-group-size-amd_32.cpp | 2 +- .../SemaSYCL/reqd-sub-group-size-amd_64.cpp | 6 +-- 3 files changed, 56 insertions(+), 4 deletions(-) create mode 100644 clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp new file mode 100644 index 0000000000000..2308f09d60eae --- /dev/null +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp @@ -0,0 +1,52 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck -check-prefix=CHECK_AMD_32 %s + +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple amdgcn-amd-amdhsa -target-cpu gfx90a -S -emit-llvm -o - %s | FileCheck -check-prefix=CHECK_AMD_64 %s + +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple nvptx-unknown-unknown -target-cpu sm_90 -S -emit-llvm -o - %s | FileCheck -check-prefix=CHECK_CUDA_32 %s + +// Check that incorrect values specified for reqd_sub_group_size are ignored. +// CDNA supports only 64 wave front size, for those GPUs allow subgroup size of +// 64. Some GPUs support both 32 and 64, for those (and the rest) only allow +// 32. For CUDA only allow 32. + +#include "sycl.hpp" + +int main() { + + sycl::queue Q; + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); + }); + + Q.submit([&](sycl::handler &h) { + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); + }); + + return 0; +} + +// CHECK_AMD_32: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_1() #0 {{.*}} +// CHECK_AMD_32-NOT: intel_reqd_sub_group_size +// CHECK_AMD_32: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_2() #0 {{.*}} !intel_reqd_sub_group_size ![[IRSGS_32:[0-9]+]] +// CHECK_AMD_32: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_3() #0 {{.*}} +// CHECK_AMD_32-NOT: intel_reqd_sub_group_size +// CHECK_AMD_32: ![[IRSGS_32]] = !{i32 32} + +// CHECK_AMD_64: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_1() #0 {{.*}} !intel_reqd_sub_group_size ![[IRSGS_64:[0-9]+]] +// CHECK_AMD_64: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_2() #0 {{.*}} +// CHECK_AMD_64-NOT: intel_reqd_sub_group_size +// CHECK_AMD_64: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_3() #0 {{.*}} +// CHECK_AMD_64-NOT: intel_reqd_sub_group_size +// CHECK_AMD_64: ![[IRSGS_64]] = !{i32 64} + +// CHECK_CUDA_32: define {{.*}} void @{{.*}}Kernel_1() #0 {{.*}} +// CHECK_CUDA_32-NOT: intel_reqd_sub_group_size +// CHECK_CUDA_32: define {{.*}} void @{{.*}}Kernel_2() #0 {{.*}} !intel_reqd_sub_group_size ![[IRSGS_32:[0-9]+]] +// CHECK_CUDA_32: define {{.*}} void @{{.*}}Kernel_3() #0 {{.*}} +// CHECK_CUDA_32-NOT: intel_reqd_sub_group_size +// CHECK_CUDA_32: ![[IRSGS_32]] = !{i32 32} diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp index 0b84abfa6b059..c97d428730dc8 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_32.cpp @@ -17,7 +17,7 @@ int main() { }); Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 32}} + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 32}} }); return 0; diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp index 6d457a8e0809d..e5694c543a77d 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-amd_64.cpp @@ -9,15 +9,15 @@ int main() { sycl::queue Q; Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); + h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); }); Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; amdgcn requires sub_group size 64}} + h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); // expected-warning {{attribute argument 32 is invalid and will be ignored; amdgcn requires sub_group size 64}} }); Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 64}} + h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); // expected-warning {{attribute argument 8 is invalid and will be ignored; amdgcn requires sub_group size 64}} }); return 0; From 82f3c5e22c72458f2da299b21fc259cfa8560055 Mon Sep 17 00:00:00 2001 From: Jakub Chlanda Date: Tue, 21 Nov 2023 08:05:23 -0500 Subject: [PATCH 5/5] PR feedback --- clang/lib/Sema/SemaDeclAttr.cpp | 10 ++-- .../reqd-sub-group-size_ignored_values.cpp | 52 ------------------- 2 files changed, 3 insertions(+), 59 deletions(-) delete mode 100644 clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 1363a2d89ae60..0d587445f729d 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -4021,11 +4021,9 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, return; } auto &TI = Context.getTargetInfo(); - if (TI.getTriple().isNVPTX() && ArgVal != 32) { + if (TI.getTriple().isNVPTX() && ArgVal != 32) Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n) << ArgVal.getSExtValue() << TI.getTriple().getArchName() << 32; - return; - } if (TI.getTriple().isAMDGPU()) { const auto HasWaveFrontSize64 = TI.getTargetOpts().FeatureMap["wavefrontsize64"]; @@ -4034,15 +4032,13 @@ void Sema::AddIntelReqdSubGroupSize(Decl *D, const AttributeCommonInfo &CI, // CDNA supports only 64 wave front size, for those GPUs allow subgroup // size of 64. Some GPUs support both 32 and 64, for those (and the rest) - // only allow 32. + // only allow 32. Warn on incompatible sizes. const auto SupportedWaveFrontSize = HasWaveFrontSize64 && !HasWaveFrontSize32 ? 64 : 32; - if (ArgVal != SupportedWaveFrontSize) { + if (ArgVal != SupportedWaveFrontSize) Diag(E->getExprLoc(), diag::warn_reqd_sub_group_attribute_n) << ArgVal.getSExtValue() << TI.getTriple().getArchName() << SupportedWaveFrontSize; - return; - } } // Check to see if there's a duplicate attribute with different values diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp deleted file mode 100644 index 2308f09d60eae..0000000000000 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size_ignored_values.cpp +++ /dev/null @@ -1,52 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple amdgcn-amd-amdhsa -target-cpu gfx1010 -S -emit-llvm -o - %s | FileCheck -check-prefix=CHECK_AMD_32 %s - -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple amdgcn-amd-amdhsa -target-cpu gfx90a -S -emit-llvm -o - %s | FileCheck -check-prefix=CHECK_AMD_64 %s - -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple nvptx-unknown-unknown -target-cpu sm_90 -S -emit-llvm -o - %s | FileCheck -check-prefix=CHECK_CUDA_32 %s - -// Check that incorrect values specified for reqd_sub_group_size are ignored. -// CDNA supports only 64 wave front size, for those GPUs allow subgroup size of -// 64. Some GPUs support both 32 and 64, for those (and the rest) only allow -// 32. For CUDA only allow 32. - -#include "sycl.hpp" - -int main() { - - sycl::queue Q; - - Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(64)]] {}); - }); - - Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(32)]] {}); - }); - - Q.submit([&](sycl::handler &h) { - h.single_task([=] [[sycl::reqd_sub_group_size(8)]] {}); - }); - - return 0; -} - -// CHECK_AMD_32: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_1() #0 {{.*}} -// CHECK_AMD_32-NOT: intel_reqd_sub_group_size -// CHECK_AMD_32: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_2() #0 {{.*}} !intel_reqd_sub_group_size ![[IRSGS_32:[0-9]+]] -// CHECK_AMD_32: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_3() #0 {{.*}} -// CHECK_AMD_32-NOT: intel_reqd_sub_group_size -// CHECK_AMD_32: ![[IRSGS_32]] = !{i32 32} - -// CHECK_AMD_64: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_1() #0 {{.*}} !intel_reqd_sub_group_size ![[IRSGS_64:[0-9]+]] -// CHECK_AMD_64: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_2() #0 {{.*}} -// CHECK_AMD_64-NOT: intel_reqd_sub_group_size -// CHECK_AMD_64: define {{.*}}amdgpu_kernel void @{{.*}}Kernel_3() #0 {{.*}} -// CHECK_AMD_64-NOT: intel_reqd_sub_group_size -// CHECK_AMD_64: ![[IRSGS_64]] = !{i32 64} - -// CHECK_CUDA_32: define {{.*}} void @{{.*}}Kernel_1() #0 {{.*}} -// CHECK_CUDA_32-NOT: intel_reqd_sub_group_size -// CHECK_CUDA_32: define {{.*}} void @{{.*}}Kernel_2() #0 {{.*}} !intel_reqd_sub_group_size ![[IRSGS_32:[0-9]+]] -// CHECK_CUDA_32: define {{.*}} void @{{.*}}Kernel_3() #0 {{.*}} -// CHECK_CUDA_32-NOT: intel_reqd_sub_group_size -// CHECK_CUDA_32: ![[IRSGS_32]] = !{i32 32}