From b99d43f1563348057682afab1dfafb5615986f86 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 27 May 2021 09:48:08 -0700 Subject: [PATCH 1/6] [SYCL] feature In SYCL 1.2.1 spec, the attributes get propagated from device functions to a kernel. The SYCL 2020 requirement mandating the avoidance of the propagation of all kernel attributes to the caller when used on a function. Attributes that should not be propagated from device functions to a kernel to match with new SYCL 2020 spec. 1. scheduler_target_fmax_mhz 2. kernel_args_restrict 3. no_global_work_offset 4. max-work-group-size 5. max-global-work-dim 6. num-simd-work-items 7. reqd-sub-group-size 8. reqd-work-group-size 9. named_sub_group_size 10. sycl_explicit_simd This patch i. keeps the SYCL 1.2.1 spec functionality and propagates the attributes with the older SYCL mode(-sycl-std=2017) ii. adds diagnostic for ignored attribute for attribute spelling [[intel:: named_sub_group_size()]] with earlier version of SYCL mode (-sycl-std=2017) since this attribute follows the SYCL 2020 Attribute Rules. iii. adds or updates tests to validate the propagating behavior with SYCL 2020 and SYCL 2017 modes. Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 53 +++++++++++-- clang/lib/Sema/SemaDeclAttr.cpp | 10 +++ clang/lib/Sema/SemaSYCL.cpp | 52 ++++++++----- .../intel-fpga-no-global-work-offset.cpp | 21 ++++- .../CodeGenSYCL/intel-max-global-work-dim.cpp | 22 +++++- .../CodeGenSYCL/intel-max-work-group-size.cpp | 22 +++++- clang/test/CodeGenSYCL/loop_fuse_device.cpp | 2 +- .../test/CodeGenSYCL/loop_fuse_ind_device.cpp | 2 +- .../test/CodeGenSYCL/num-simd-work-items.cpp | 23 +++++- .../test/CodeGenSYCL/reqd-sub-group-size.cpp | 31 ++++++-- .../test/CodeGenSYCL/reqd-work-group-size.cpp | 26 ++++++- .../CodeGenSYCL/scheduler-target-fmax-mhz.cpp | 13 +++- .../CodeGenSYCL/sycl-multi-kernel-attr.cpp | 23 +++++- .../check-notdirect-attribute-propagation.cpp | 6 +- .../intel-fpga-no-global-work-offset.cpp | 26 ++++++- .../intel-max-global-work-dim-device.cpp | 16 +++- .../intel-max-work-group-size-device.cpp | 21 ++++- .../intel-reqd-work-group-size-device.cpp | 77 +++++++++++++++++-- clang/test/SemaSYCL/intel-restrict.cpp | 18 ++++- clang/test/SemaSYCL/loop_fusion_ast.cpp | 2 +- .../SemaSYCL/num_simd_work_items_device.cpp | 23 +++++- .../redeclaration-attribute-propagation.cpp | 6 +- .../SemaSYCL/reqd-sub-group-size-device.cpp | 43 ++++++++++- .../SemaSYCL/reqd-work-group-size-device.cpp | 50 +++++++++--- .../scheduler_target_fmax_mhz_ast.cpp | 18 ++++- clang/test/SemaSYCL/sub-group-size-ast.cpp | 75 ++++++++++++++++++ clang/test/SemaSYCL/sycl-esimd-ast.cpp | 55 +++++++++++++ clang/test/SemaSYCL/sycl-esimd.cpp | 2 +- 28 files changed, 659 insertions(+), 79 deletions(-) create mode 100644 clang/test/SemaSYCL/sub-group-size-ast.cpp create mode 100644 clang/test/SemaSYCL/sycl-esimd-ast.cpp diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index e5eb377cbfe07..f8e3608720943 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -415,6 +415,12 @@ def SYCLSimdDocs : Documentation { The compiler may decide to compile such functions using different optimization and code generation pipeline. Also, this attribute is used to distinguish ESIMD private globals from regular SYCL global variables. + + When conforming to SYCL 1.2.1 extension specifications, when called from a + device kernel, the ``intel::sycl_explicit_simd`` is propagated from the + function it is applied to, to the kernel. When conforming to SYCL 2020 + extension specifications, the attribute is not propagated to the kernel. + }]; } @@ -2443,8 +2449,10 @@ lambda capture, or function object member, of the callable to which the attribute was applied. This effect is equivalent to annotating restrict on **all** kernel pointer arguments in an OpenCL or SPIR-V kernel. -If ``intel::kernel_args_restrict`` is applied to a function called from a device -kernel, the attribute is not ignored and it is propagated to the kernel. +When conforming to SYCL 1.2.1 extension specifications, when called from a +device kernel, the ``intel::kernel_args_restrict`` is propagated from the +function it is applied to, to the kernel. When conforming to SYCL 2020 +extension specifications, the attribute is not propagated to the kernel. The attribute forms an unchecked assertion, in that implementations do not need to check/confirm the pre-condition in any way. If a user applies @@ -2482,8 +2490,11 @@ def SYCLIntelNumSimdWorkItemsAttrDocs : Documentation { let Content = [{ Applies to a device function/lambda function. Indicates the number of work items that should be processed in parallel. Valid values are positive integers. -If ``intel::num_simd_work_items`` is applied to a function called from a -device kernel, the attribute is not ignored and it is propagated to the kernel. + +When conforming to SYCL 1.2.1 extension specifications, when called from a +device kernel, the ``intel::num_simd_work_items`` is propagated from the +function it is applied to, to the kernel. When conforming to SYCL 2020 +extension specifications, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -2633,6 +2644,11 @@ allows the Y and Z arguments to be optional. If not provided by the user, the value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more details. +When conforming to SYCL 1.2.1 specifications, when called from a device kernel, +the ``intel::reqd_work_group_size`` or [[sycl::reqd_work_group_size() is +propagated from the function it is applied to, to the kernel. When conforming +to SYCL 2020 specifications, the attribute is not propagated to the kernel. + .. code-block:: c++ [[sycl::reqd_work_group_size(4, 4, 4)]] void foo() {} @@ -2773,8 +2789,11 @@ Applies to a device function/lambda function. Indicates the maximum dimensions of a work group. Values must be positive integers. This is similar to reqd_work_group_size, but allows work groups that are smaller or equal to the specified sizes. -If ``intel::max_work_group_size`` is applied to a function called from a -device kernel, the attribute is not ignored and it is propagated to the kernel. + +When conforming to SYCL 1.2.1 specifications, when called from a device kernel, +the ``intel::max_work_group_size`` is propagated from the function it is +applied to, to the kernel. When conforming to SYCL 2020 specifications, the +attribute is not propagated to the kernel. .. code-block:: c++ @@ -2805,8 +2824,11 @@ Applies to a device function/lambda function or function call operator (of a function object). Indicates the largest valid global work dimension that will be accepted when running the kernel on a device. Valid values are integers in a range of [0, 3]. -If ``intel::max_global_work_dim`` is applied to a function called from a -device kernel, the attribute is not ignored and it is propagated to the kernel. + +When conforming to SYCL 1.2.1 extension specifications, when called from a +device kernel, the ``intel::max_global_work_dim`` is propagated from the +function it is applied to, to the kernel. When conforming to SYCL 2020 +extension specifications, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -2863,6 +2885,11 @@ device operation, guiding the FPGA backend to insert the appropriate number of registers to break-up the combinational logic circuit, and thereby controlling the length of the longest combinational path. +When conforming to SYCL 1.2.1 extension specifications, when called from a +device kernel, the ``intel::scheduler_target_fmax_mhz`` is propagated from the +function it is applied to, to the kernel. When conforming to SYCL 2020 +extension specifications, the attribute is not propagated to the kernel. + .. code-block:: c++ [[intel::scheduler_target_fmax_mhz(4)]] void foo() {} @@ -2893,6 +2920,11 @@ function object). If 1, compiler doesn't use the global work offset values for the device function. Valid values are 0 and 1. If used without argument, value of 1 is set implicitly. +When conforming to SYCL 1.2.1 extension specifications, when called from a +device kernel, the ``intel::no_global_work_offset`` is propagated from the +function it is applied to, to the kernel. When conforming to SYCL 2020 +extension specifications, the attribute is not propagated to the kernel. + .. code-block:: c++ [[intel::no_global_work_offset]] @@ -4567,6 +4599,11 @@ In addition to device functions, the required sub-group size attribute may also be specified in the definition of a named functor object and lambda functions, as in the examples below: +When conforming to SYCL 1.2.1 specifications, when called from a device kernel, +the ``intel::reqd_sub_group_size`` is propagated from the function it is +applied to, to the kernel. When conforming to SYCL 2020 specifications, the +attribute is not propagated to the kernel. + .. code-block:: c++ class Functor diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index f1edc1ddff801..7c32b05c85b02 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -328,6 +328,14 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, return; } + // Diagnose the [[intel::named_sub_group_size]] attribute spelling + // in earlier SYCL mode (SYCL 2017) as being an ignored attribute. + if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2017 && + A.getKind() == ParsedAttr::AT_IntelNamedSubGroupSize) { + Diag(A.getLoc(), diag::warn_attribute_ignored) << A; + return; + } + // Diagnose SYCL 2017 spellings in later SYCL modes. if (LangOpts.getSYCLVersion() > LangOptions::SYCL_2017) { // All attributes in the cl vendor namespace are deprecated in favor of a @@ -3361,6 +3369,8 @@ Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D, static void handleIntelNamedSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { + S.CheckDeprecatedSYCLAttributeSpelling(AL); + StringRef SizeStr; SourceLocation Loc; if (AL.isArgIdent(0)) { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c2c11765a8bc2..76ccb30e09c46 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -343,15 +343,40 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, if (!FD->hasAttrs()) return; - llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { - // FIXME: Make this list self-adapt as new SYCL attributes are added. - return isa(A); - }); + // Attributes that should be propagated from device functions to a kernel + // in SYCL 1.2.1. + if (S.getASTContext().getLangOpts().getSYCLVersion() < + LangOptions::SYCL_2020) { + llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { + // FIXME: Make this list self-adapt as new SYCL attributes are added. + return isa(A); + }); + // Attributes that should not be propagated from device functions to a + // kernel in SYCL 1.2.1. + if (DirectlyCalled) { + llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), + [](Attr *A) { return isa(A); }); + } + } else { + // Attributes that should not be propagated from device functions to a + // kernel in SYCL 2020. + if (DirectlyCalled) { + llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { + return isa< + SYCLIntelFPGAMaxConcurrencyAttr, + SYCLIntelFPGADisableLoopPipeliningAttr, SYCLSimdAttr, + SYCLIntelKernelArgsRestrictAttr, ReqdWorkGroupSizeAttr, + SYCLIntelNumSimdWorkItemsAttr, SYCLIntelSchedulerTargetFmaxMhzAttr, + SYCLIntelNoGlobalWorkOffsetAttr, SYCLIntelMaxWorkGroupSizeAttr, + IntelReqdSubGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr, + IntelNamedSubGroupSizeAttr, SYCLIntelFPGAInitiationIntervalAttr>(A); + }); + } + } // Allow the kernel attribute "use_stall_enable_clusters" only on lambda // functions and function objects called directly from a kernel. @@ -366,15 +391,6 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, FD->dropAttr(); } } - - // Attributes that should not be propagated from device functions to a kernel. - if (DirectlyCalled) { - llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { - return isa(A); - }); - } } class DiagDeviceFunction : public RecursiveASTVisitor { diff --git a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp index f8edf692f66e2..0ca09d453b28d 100644 --- a/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/CodeGenSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -19,6 +20,8 @@ class Functor { template [[intel::no_global_work_offset(N)]] void func() {} +[[intel::no_global_work_offset(1)]] void func1() {} + int main() { q.submit([&](handler &h) { Foo boo; @@ -30,12 +33,26 @@ int main() { h.single_task( []() [[intel::no_global_work_offset(0)]]{}); + // Test class template argument. Functor<1> f; h.single_task(f); +#if defined(SYCL2017) + // Test template argument with propagated function attribute. h.single_task([]() { func<1>(); }); + + // Test attribute is propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif //SYCL2020 }); return 0; } @@ -45,5 +62,7 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} ![[NUM4:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !no_global_work_offset ![[NUM5]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !no_global_work_offset ![[NUM5]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} !no_global_work_offset ![[NUM5]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7"() #0 {{.*}} ![[NUM5]] // CHECK-NOT: ![[NUM4]] = !{i32 0} // CHECK: ![[NUM5]] = !{} diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp index c127e98280114..09c4ccd3d0fcd 100644 --- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -19,6 +20,8 @@ class Functor { template [[intel::max_global_work_dim(N)]] void func() {} +[[intel::max_global_work_dim(2)]] void func1() {} + int main() { q.submit([&](handler &h) { Foo boo; @@ -27,12 +30,26 @@ int main() { h.single_task( []() [[intel::max_global_work_dim(2)]]{}); + // Test class template argument. Functor<2> f; h.single_task(f); +#if defined(SYCL2017) + // Test template argument with propagated function attribute. h.single_task([]() { func<2>(); }); + + // Test attribute is propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif //SYCL2020 }); return 0; } @@ -41,5 +58,8 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !max_global_work_dim ![[NUM2:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !max_global_work_dim ![[NUM2]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !max_global_work_dim ![[NUM2]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !max_global_work_dim ![[NUM2]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} ![[NUM0:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM2]] = !{i32 2} +// CHECK: ![[NUM0]] = !{} diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp index f59c8e91f9191..fd9acd1de5b7f 100644 --- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -24,6 +25,8 @@ class Functor { template [[intel::max_work_group_size(N, N1, N2)]] void func() {} +[[intel::max_work_group_size(10, 10, 10)]] void func1() {} + int main() { q.submit([&](handler &h) { Foo boo; @@ -35,12 +38,26 @@ int main() { Bar bar; h.single_task(bar); + // Test class template argument. Functor<2, 2, 2> f; h.single_task(f); +#if defined(SYCL2017) + // Test template argument with propagated function attribute. h.single_task([]() { func<4, 4, 4>(); }); + + // Test attribute is propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif //SYCL2020 }); return 0; } @@ -50,8 +67,11 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !max_work_group_size ![[NUM6:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !max_work_group_size ![[NUM2:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !max_work_group_size ![[NUM4:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} !max_work_group_size ![[NUM10:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7"() #0 {{.*}} ![[NUM0:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1, i32 1, i32 1} // CHECK: ![[NUM8]] = !{i32 8, i32 8, i32 8} // CHECK: ![[NUM6]] = !{i32 6, i32 3, i32 1} // CHECK: ![[NUM2]] = !{i32 2, i32 2, i32 2} // CHECK: ![[NUM4]] = !{i32 4, i32 4, i32 4} +// CHECK: ![[NUM0]] = !{} diff --git a/clang/test/CodeGenSYCL/loop_fuse_device.cpp b/clang/test/CodeGenSYCL/loop_fuse_device.cpp index 645bb908d24c1..9e1e92fc36260 100644 --- a/clang/test/CodeGenSYCL/loop_fuse_device.cpp +++ b/clang/test/CodeGenSYCL/loop_fuse_device.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp b/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp index 49cc7640e1de7..ca3eb4445e1bd 100644 --- a/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp +++ b/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp index a5c28285f0ef0..d33130fb7bf6c 100644 --- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp +++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -19,6 +20,8 @@ class Functor { template [[intel::num_simd_work_items(N)]] void func() {} +[[intel::num_simd_work_items(10)]] void func1() {} + int main() { q.submit([&](handler &h) { Foo boo; @@ -27,12 +30,26 @@ int main() { h.single_task( []() [[intel::num_simd_work_items(42)]]{}); + // Test class template argument. Functor<2> f; h.single_task(f); + #if defined(SYCL2017) + // Test template argument with propagated function attribute. h.single_task([]() { func<4>(); }); + + // Test attribute is propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2020 }); return 0; } @@ -41,7 +58,11 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2"() #0 {{.*}} !num_simd_work_items ![[NUM42:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !num_simd_work_items ![[NUM2:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !num_simd_work_items ![[NUM4:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !num_simd_work_items ![[NUM10:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} ![[NUM0:[0-9]+]] // CHECK: ![[NUM1]] = !{i32 1} // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM2]] = !{i32 2} // CHECK: ![[NUM4]] = !{i32 4} +// CHECK: ![[NUM10]] = !{i32 10} +// CHECK: ![[NUM0]] = !{} diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp index aa9c104cfc61f..eac1603ccf1d4 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -28,6 +29,8 @@ class Functor2 { template [[intel::reqd_sub_group_size(N)]] void func() {} +[[intel::reqd_sub_group_size(10)]] void func1() {} + int main() { q.submit([&](handler &h) { Functor16 f16; @@ -39,12 +42,26 @@ int main() { h.single_task( []() [[intel::reqd_sub_group_size(4)]]{}); +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2020 + + // Test class template argument. Functor2<2> f2; - h.single_task(f2); + h.single_task(f2); - h.single_task([]() { +#if defined(SYCL2017) + // Test template argument with propagated function attribute. + h.single_task([]() { func<2>(); }); + + // Test attribute is propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2017 }); return 0; } @@ -52,9 +69,13 @@ int main() { // 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_name4"() #0 {{.*}} ![[SGSIZE0:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE2]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7"() #0 {{.*}} !intel_reqd_sub_group_size ![[SGSIZE10:[0-9]+]] // CHECK: ![[SGSIZE16]] = !{i32 16} // CHECK: ![[SGSIZE8]] = !{i32 8} // CHECK: ![[SGSIZE4]] = !{i32 4} +// CHECK: ![[SGSIZE0]] = !{} // CHECK: ![[SGSIZE2]] = !{i32 2} +// CHECK: ![[SGSIZE10]] = !{i32 10} diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index 11047dbff3168..ddc4fde8f1078 100644 --- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -28,6 +29,8 @@ class FunctorTemp { template [[cl::reqd_work_group_size(N, N1, N2)]] void func() {} +[[cl::reqd_work_group_size(10, 10, 10)]] void func1() {} + int main() { q.submit([&](handler &h) { Functor32x16x16 f32x16x16; @@ -39,15 +42,30 @@ int main() { h.single_task( []() [[cl::reqd_work_group_size(8, 8, 8)]]{}); + // Test class template argument. FunctorTemp<2, 2, 2> ft; h.single_task(ft); +#if defined(SYCL2017) + // Test template argument with propagated function attribute. h.single_task([]() { func<8, 4, 4>(); }); + // Test attribute is propagated. h.single_task( + []() { func1(); }); +#endif // SYCL2017 + + // Test attribute is applied on lambda. + h.single_task( []() [[cl::reqd_work_group_size(1, 8, 2)]]{}); + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL2020 }); return 0; } @@ -57,10 +75,14 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE88:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE22:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE44:[0-9]+]] -// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE10:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE2:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8"() #0 {{.*}} ![[NUM0:[0-9]+]] // CHECK: ![[WGSIZE32]] = !{i32 16, i32 16, i32 32} // CHECK: ![[WGSIZE8]] = !{i32 1, i32 1, i32 8} // CHECK: ![[WGSIZE88]] = !{i32 8, i32 8, i32 8} // CHECK: ![[WGSIZE22]] = !{i32 2, i32 2, i32 2} // CHECK: ![[WGSIZE44]] = !{i32 4, i32 4, i32 8} +// CHECK: ![[WGSIZE10]] = !{i32 10, i32 10, i32 10} // CHECK: ![[WGSIZE2]] = !{i32 2, i32 8, i32 1} +// CHECK: ![[NUM0]] = !{} diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp index 715e2dd204459..cf28dd18caa46 100644 --- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp +++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -35,6 +36,7 @@ int main() { Functor<7> f; h.single_task(f); +#if defined(SYCL2017) // Test attribute is propagated. h.single_task( []() { bar(); }); @@ -42,6 +44,13 @@ int main() { // Test function template argument. h.single_task( []() { zoo<75>(); }); +#endif //SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { bar(); }); +#endif //SYCL2020 }); return 0; } @@ -51,8 +60,10 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM7:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]] // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM75:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} ![[NUM0:[0-9]+]] // CHECK: ![[NUM5]] = !{i32 5} // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM7]] = !{i32 7} // CHECK: ![[NUM2]] = !{i32 2} // CHECK: ![[NUM75]] = !{i32 75} +// CHECK: ![[NUM0]] = !{} diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index f2cecd45fcf8d..49c2ec5177dc1 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" @@ -24,6 +25,8 @@ class Functor2 { template [[intel::reqd_work_group_size(N, N1, N2)]] void func() {} +[[intel::reqd_work_group_size(10, 10, 10)]] void func1() {} + int main() { q.submit([&](handler &h) { Functor foo; @@ -32,12 +35,26 @@ int main() { Functor1 foo1; h.single_task(foo1); + // Test clss template argument. Functor2<2, 2, 2> foo2; h.single_task(foo2); +#if defined(SYCL2017) + // Test template argument with propagated function attribute. h.single_task([]() { func<8, 4, 4>(); }); + + // Test attribute is propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL 2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task( + []() { func1(); }); +#endif // SYCL 2020 }); return 0; } @@ -46,9 +63,13 @@ int main() { // 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_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: define {{.*}}spir_kernel void @{{.*}}kernel_name5"() #0 {{.*}} !reqd_work_group_size ![[WGSIZE10:[0-9]+]] +// CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6"() #0 {{.*}} ![[NUM0:[0-9]+]] // CHECK: ![[WGSIZE]] = !{i32 16, i32 16, i32 32} // CHECK: ![[SGSIZE]] = !{i32 4} // CHECK: ![[WGSIZE1]] = !{i32 32, i32 32, i32 64} // CHECK: ![[SGSIZE1]] = !{i32 2} // CHECK: ![[WGSIZE2]] = !{i32 2, i32 2, i32 2} // CHECK: ![[WGSIZE3]] = !{i32 4, i32 4, i32 8} +// CHECK: ![[WGSIZE10]] = !{i32 10, i32 10, i32 10} +// CHECK: ![[NUM0]] = !{} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index a7bab5a2f43c1..03d605080cdd7 100644 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -triple spir64 -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s #ifndef TRIGGER_ERROR [[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index b3c2d4b8e40c2..e342d861385ee 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -sycl-std=2017 -ast-dump -DSYCL2017 -verify -pedantic %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -sycl-std=2020 -ast-dump -DSYCL2020 -verify -pedantic %s #include "sycl.hpp" @@ -11,6 +12,8 @@ struct FuncObj { [[intelfpga::no_global_work_offset]] void operator()() const {} }; +[[intel::no_global_work_offset(1)]] void func() {} + int main() { q.submit([&](handler &h) { // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} @@ -61,7 +64,26 @@ int main() { h.single_task( []() [[intel::no_global_work_offset(0), // expected-note {{previous attribute is here}} - intel::no_global_work_offset(1)]]{}); // expected-warning{{attribute 'no_global_work_offset' is already applied with different arguments}} + intel::no_global_work_offset(1)]]{}); // expected-warning{{attribute 'no_global_work_offset' is already applied with different arguments}} + +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9 + // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + h.single_task( + []() { func(); }); +#endif // SYCL2020 + +#if defined(SYCL2017) + // Test attribute is propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kerne10 + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task( + []() { func(); }); +#endif // SYCL2017 }); return 0; } diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp index c07c12b5d98d8..82270448d3332 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -1,5 +1,7 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s #include "sycl.hpp" @@ -75,7 +77,16 @@ int main() { // expected-note@+2 {{did you mean to use 'intel::max_global_work_dim' instead?}} h.single_task( []() [[intelfpga::max_global_work_dim(2)]]{}); +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 + // CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + h.single_task( + []() { func_do_not_ignore(); }); +#endif // SYCL2020 +#if defined(SYCL2017) + // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -137,6 +148,7 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +#endif // SYCL2017 #ifdef TRIGGER_ERROR [[intel::max_global_work_dim(1)]] int Var = 0; // expected-error{{'max_global_work_dim' attribute only applies to functions}} diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp index f190e5920ee38..ea7e2718ab0c0 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp @@ -1,5 +1,7 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s #include "sycl.hpp" @@ -34,6 +36,7 @@ struct Func { }; #ifdef TRIGGER_ERROR +#if defined(SYCL2020) struct DAFuncObj { [[intel::max_work_group_size(4, 4, 4)]] [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \ @@ -41,6 +44,7 @@ struct DAFuncObj { // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}} void operator()() const {} }; +#endif // SYCL2020 #endif // TRIGGER_ERROR int main() { @@ -74,6 +78,16 @@ int main() { h.single_task( []() [[intelfpga::max_work_group_size(8, 8, 8)]]{}); +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 + // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + h.single_task( + []() { func_do_not_ignore(); }); +#endif // SYCL2020 + +#if defined(SYCL2017) + // Test attribute is propagated // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -87,6 +101,7 @@ int main() { // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} h.single_task( []() { func_do_not_ignore(); }); +#endif // SYCL2017 // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} @@ -133,8 +148,10 @@ int main() { []() [[intel::max_work_group_size(16, 16, 16), // expected-note{{conflicting attribute is here}} intel::max_work_group_size(2, 2, 2)]]{}); // expected-warning{{attribute 'max_work_group_size' is already applied with different arguments}} +#if defined(SYCL2020) h.single_task( DAFuncObj()); +#endif // SYCL2020 #endif // TRIGGER_ERROR }); diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp index 74449a3777fa8..a0047824dbf33 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -1,5 +1,7 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -fsyntax-only -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -DSYCL2017 -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -DSYCL2020 -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s #include "sycl.hpp" @@ -21,25 +23,38 @@ void bar() { } #else +#if defined(SYCL2017) [[intel::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} - [[intel::reqd_work_group_size(16)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} - [[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} +#endif // SYCL2017 + +#if defined(SYCL2020) +[[intel::reqd_work_group_size(4)]] void f4x1x1() {} // OK +[[intel::reqd_work_group_size(32)]] void f32x1x1() {} // OK +[[intel::reqd_work_group_size(16)]] void f16x1x1() {} // OK +[[intel::reqd_work_group_size(16, 16)]] void f16x16x1() {} // OK +[[intel::reqd_work_group_size(32, 32)]] void f32x32x1() {} // OK +[[intel::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // OK +#endif // SYCL2020 #ifdef TRIGGER_ERROR +#if defined(SYCL2020) class Functor32 { public: [[cl::reqd_work_group_size(32)]] void operator()() const {} // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \ // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} }; +#endif // SYCL2020 #endif // TRIGGER_ERROR +[[intel::reqd_work_group_size(16, 16, 16)]] void func() {} + class Functor33 { public: // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} @@ -67,12 +82,14 @@ class Functor16x16x16 { [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; +#if defined(SYCL2017) class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: [[intel::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} f4x1x1(); } }; +#endif // SYCL2017 class Functor { public: @@ -80,26 +97,32 @@ class Functor { f4x1x1(); } }; - +#if defined(SYCL2020) class FunctorAttr { public: __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \ // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} }; +#endif // SYCL2020 int main() { q.submit([&](handler &h) { Functor16 f16; h.single_task(f16); +#if defined(SYCL2017) + // Test attribute is propagated. Functor f; h.single_task(f); +#endif // SYCL2017 Functor16x16x16 f16x16x16; h.single_task(f16x16x16); +#if defined(SYCL2020) FunctorAttr fattr; h.single_task(fattr); +#endif // SYCL2020 Functor33 f33; h.single_task(f33); @@ -111,6 +134,8 @@ int main() { f32x32x32(); }); #ifdef TRIGGER_ERROR +#if defined(SYCL2017) + // Test attribute is propagated. Functor8 f8; h.single_task(f8); @@ -128,13 +153,55 @@ int main() { f32x32x32(); f32x32x1(); }); +#endif // SYCL2017 +#if defined(SYCL2020) + // Test attribute is not propagated. + h.single_task([]() { // OK + f4x1x1(); + f32x1x1(); + }); + h.single_task([]() { // OK + f16x1x1(); + f16x16x1(); + }); + + h.single_task([]() { // OK + f32x32x32(); + f32x32x1(); + }); +#endif // SYCL2020 // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} h.single_task([[intel::reqd_work_group_size(32, 32, 32)]][]() { f32x32x32(); }); #endif // TRIGGER_ERROR + +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}class kernel_name13 + // CHECK-NOT: ReqdWorkGroupSizeAttr {{.*}} + h.single_task( + []() { func(); }); +#endif // SYCL2020 + +#if defined(SYCL2017) + // Test attribute is propagated. + // CHECK-LABEL: FunctionDecl {{.*}}class kernel_name14 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + h.single_task( + []() { func(); }); +#endif // SYCL2017 }); return 0; } diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index bd51eb6b5c055..2970d601f2c02 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -1,5 +1,8 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -Wno-sycl-2017-compat -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DCHECKDIAG -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DCHECKDIAG -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s + [[intel::kernel_args_restrict]] void func_do_not_ignore() {} @@ -26,8 +29,19 @@ int main() { kernel( []() [[intel::kernel_args_restrict]] {}); +#if defined(SYCL2017) + // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelKernelArgsRestrictAttr kernel( []() { func_do_not_ignore(); }); +#endif // SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr + kernel( + []() { func_do_not_ignore(); }); +#endif // SYCL2020 } diff --git a/clang/test/SemaSYCL/loop_fusion_ast.cpp b/clang/test/SemaSYCL/loop_fusion_ast.cpp index 72407b70b65eb..20bf4f30bca82 100644 --- a/clang/test/SemaSYCL/loop_fusion_ast.cpp +++ b/clang/test/SemaSYCL/loop_fusion_ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s // Tests for AST of Intel FPGA loop fusion function attributes #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/num_simd_work_items_device.cpp b/clang/test/SemaSYCL/num_simd_work_items_device.cpp index bfa01d8b782cf..5e88596711dc9 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -1,5 +1,7 @@ -// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -Wno-sycl-2017-compat -ast-dump | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s #include "sycl.hpp" @@ -134,18 +136,22 @@ struct TRIFuncObjBad8 { [[intel::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} [[intel::num_simd_work_items(2)]] void func2(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +#if defined(SYCL2020) [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} [[cl::reqd_work_group_size(4, 2, 3)]] void func3(); // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} [[cl::reqd_work_group_size(4, 2, 3)]] // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} [[intel::num_simd_work_items(2)]] void func4(); // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} +#endif // SYCL2020 // If the declaration has a __attribute__((reqd_work_group_size())) // attribute, tests that check if the work group size attribute argument // (the last argument) can be evenly divided by the [[intel::num_simd_work_items()]] // attribute. +#if defined(SYCL2020) [[intel::num_simd_work_items(2)]] // expected-error{{'num_simd_work_items' attribute must evenly divide the work-group size for the 'reqd_work_group_size' attribute}} __attribute__((reqd_work_group_size(4, 2, 5))) void func5(); // expected-note{{conflicting attribute is here}} expected-warning {{attribute 'reqd_work_group_size' is deprecated}} expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} +#endif // SYCL2020 // Tests for incorrect argument values for Intel FPGA num_simd_work_items and reqd_work_group_size function attributes struct TRIFuncObjBad9 { @@ -238,9 +244,11 @@ struct TRIFuncObjGood4 { operator()() const {} }; +#if defined(SYCL2020) [[intel::num_simd_work_items(2)]] __attribute__((reqd_work_group_size(3, 2, 6))) void func6(); // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \ expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} +#endif // SYCL2020 int main() { q.submit([&](handler &h) { @@ -261,6 +269,16 @@ int main() { h.single_task( []() [[intelfpga::num_simd_work_items(8)]]{}); +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 + // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} + h.single_task( + []() { func_do_not_ignore(); }); +#endif // SYCL2020 + +#if defined(SYCL2017) + // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -336,6 +354,7 @@ int main() { // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 4 // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} +#endif // SYCL2017 #ifdef TRIGGER_ERROR [[intel::num_simd_work_items(0)]] int Var = 0; // expected-error{{'num_simd_work_items' attribute only applies to functions}} diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp index 448c36d5fdc98..56b70a64b3b7d 100644 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp @@ -1,6 +1,6 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -Wno-sycl-2017-compat | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -sycl-std=2017 -Wno-sycl-2017-compat -verify +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -sycl-std=2017 -DTRIGGER_ERROR -Wno-sycl-2017-compat -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 -Wno-sycl-2017-compat | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp index d4694411f0fb8..36a269a456f1d 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp @@ -1,14 +1,25 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -DTRIGGER_ERROR -DSYCL2017 -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -DTRIGGER_ERROR -DSYCL2020 -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -DSYCL2020 %s #include "sycl.hpp" using namespace cl::sycl; queue q; +[[intel::reqd_sub_group_size(16)]] void func() {} + +#if defined(SYCL2017) [[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}} +#endif // SYCL2017 + +#if defined(SYCL2020) +[[intel::reqd_sub_group_size(4)]] void foo() {} // OK +[[intel::reqd_sub_group_size(32)]] void baz() {} // OK +#endif // SYCL2020 // No diagnostic is emitted because the arguments match. [[intel::reqd_sub_group_size(12)]] void bar(); @@ -32,6 +43,7 @@ class Functor16 { [[intel::reqd_sub_group_size(16)]] void operator()() const {} }; +#if defined(SYCL2017) class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} @@ -39,11 +51,13 @@ class Functor8 { // expected-error {{conflicting attributes applied to a SYCL ke } }; +#endif // SYCL2017 class Functor4 { public: [[intel::reqd_sub_group_size(12)]] void operator()() const {} }; +#if defined(SYCL2017) class Functor { public: void operator()() const { @@ -51,11 +65,15 @@ class Functor { } }; +#endif // SYCL2017 + int main() { q.submit([&](handler &h) { Functor16 f16; h.single_task(f16); +#if defined(SYCL2017) + // Test attribute is propagated. Functor f; h.single_task(f); @@ -67,7 +85,8 @@ int main() { foo(); baz(); }); -#endif +#endif // TRIGGER_ERROR +#endif // SYCL2017 h.single_task([]() [[intel::reqd_sub_group_size(2)]]{}); h.single_task([]() [[intel::reqd_sub_group_size(4)]] { foo(); }); @@ -75,6 +94,24 @@ int main() { Functor4 f4; h.single_task(f4); + +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}class kernel_name9 + // CHECK-NOT: IntelReqdSubGroupSizeAttr {{.*}} + h.single_task( + []() { func(); }); +#endif // SYCL2020 + +#if defined(SYCL2017) + // Test attribute is propagated. + // CHECK-LABEL: FunctionDecl {{.*}}class kernel_name10 + // CHECK: IntelReqdSubGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + h.single_task( + []() { func(); }); +#endif // SYCL2017 }); return 0; } diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index c737afe30ef50..1d668e8945961 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -1,20 +1,30 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -DSYCL2017 -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -DSYCL2020 -verify %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s #include "sycl.hpp" using namespace cl::sycl; queue q; -[[sycl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} +#if defined(SYCL2017) +[[sycl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} \ + // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} - [[sycl::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} - [[sycl::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} +#endif // SYCL2017 + +#if defined(SYCL2020) +[[sycl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // OK +[[sycl::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // OK +[[sycl::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // OK +[[sycl::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // OK +[[sycl::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // OK +[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // OK // No diagnostic because the attributes are synonyms with identical behavior. [[intel::reqd_work_group_size(4, 4, 4)]] void four(); @@ -67,12 +77,14 @@ class Functor32 { // expected-error@+1{{'reqd_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}} [[sycl::reqd_work_group_size(32, 1, 1)]] [[sycl::reqd_work_group_size(1, 1, 32)]] void operator()() const {} }; -#endif +#endif // TRIGGER_ERROR class Functor16x16x16 { public: [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; +#endif // SYCL2020 +#if defined(SYCL2017) class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: [[sycl::reqd_work_group_size(1, 1, 8)]] void operator()() const { // expected-note {{conflicting attribute is here}} @@ -86,21 +98,31 @@ class Functor { f4x1x1(); } }; +#endif // SYCL2017 +#if defined(SYCL2020) class FunctorAttr { public: __attribute__((reqd_work_group_size(128, 128, 128))) void operator()() const {} // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \ // expected-note {{did you mean to use '[[sycl::reqd_work_group_size]]' instead?}} }; +#endif // SYCL2020 int main() { q.submit([&](handler &h) { + +#if defined(SYCL2020) Functor16 f16; h.single_task(f16); +#endif // SYCL2020 +#if defined(SYCL2017) + // Test attribute is propagated. Functor f; h.single_task(f); +#endif // SYCL2017 +#if defined(SYCL2020) Functor16x16x16 f16x16x16; h.single_task(f16x16x16); @@ -110,14 +132,19 @@ int main() { h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32), sycl::reqd_work_group_size(32, 32, 32)]] { f32x32x32(); }); +#endif // SYCL2020 #ifdef TRIGGER_ERROR +#if defined(SYCL2017) + // Test attribute is propagated. Functor8 f8; h.single_task(f8); - +#endif // SYCL2017 +#if defined(SYCL2020) Functor32 f32; h.single_task(f32); - +#endif // SYCL2020 +#if defined(SYCL2017) h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} f4x1x1(); f32x1x1(); @@ -132,12 +159,13 @@ int main() { f32x32x32(); f32x32x1(); }); - +#endif // SYCL2017 +#if defined(SYCL2020) // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { f32x32x32(); }); - +#endif // SYCL2020 #endif }); return 0; diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp index eca66915670f3..e3f53b23f5807 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -ast-dump -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -ast-dump -DSYCL2020 %s // Tests for AST of Intel FPGA scheduler_target_fmax_mhz function attribute. #include "sycl.hpp" @@ -61,11 +62,23 @@ class KernelFunctor2 { int main() { deviceQueue.submit([&](sycl::handler &h) { +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 + // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr + KernelFunctor f1; + h.single_task(f1); +#endif // SYCL2020 + +#if defined(SYCL2017) + // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr KernelFunctor f1; h.single_task(f1); +#endif // SYCL2017 + // Test class template argument. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -84,6 +97,8 @@ int main() { h.single_task( []() [[intel::scheduler_target_fmax_mhz(4)]]{}); +#if defined(SYCL2017) + // Test function template argument. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -93,6 +108,7 @@ int main() { // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75 h.single_task( []() { func3<75>(); }); +#endif // SYCL2017 // Ignore duplicate attribute. h.single_task( diff --git a/clang/test/SemaSYCL/sub-group-size-ast.cpp b/clang/test/SemaSYCL/sub-group-size-ast.cpp new file mode 100644 index 0000000000000..5f3cb55e2bf2b --- /dev/null +++ b/clang/test/SemaSYCL/sub-group-size-ast.cpp @@ -0,0 +1,75 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -DSYCL2017 -verify %s + +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,integer -DSYCL2020 %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,primary -DSYCL2020 %s + +// Validate the semantic analysis checks for the named_sub_group_size attribute. +// The attribute can be only be applied to non-sycl-kernel/ +// non-sycl-device functions if they match the kernel they are being called +// from. + +#include "Inputs/sycl.hpp" + +#if defined(SYCL2017) +// Test that we get ignored attribute warning when using +// a [[intel::named_sub_group_size()]] attribute spelling while not +// in SYCL 2020 mode. +[[intel::named_sub_group_size(automatic)]] void func_ignore(); // expected-warning {{'named_sub_group_size' attribute ignored}} +#endif // SYCL2017 + +#if defined(SYCL2020) +// The kernel has an attribute. +void calls_kernel_1() { + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_1vE7Kernel1 + // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic + sycl::kernel_single_task([]() [[intel::named_sub_group_size(automatic)]] { + }); +} + +struct Functor { + [[intel::named_sub_group_size(automatic)]] void operator()() const { + } +}; + +struct Functor1 { + [[intel::named_sub_group_size(primary)]] void operator()() const { + } +}; + +// Test attributes get propgated to the kernel. +void calls_kernel_2() { + Functor F; + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_2vE7Kernel2 + // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic + sycl::kernel_single_task(F); + + Functor1 F1; + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_3vE7Kernel3 + // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Primary + sycl::kernel_single_task(F1); +} + +// Test ttribute does not get propgated to the kernel. +[[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc + +void calls_kernel_3() { + // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_4vE7Kernel4 + // CHECK-NOT : IntelNamedSubGroupSizeAttr {{.*}} + sycl::kernel_single_task([]() { // #Kernel4 + // primary-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}} + // primary-note@#Kernel4{{kernel declared here}} + AttrFunc(); + }); +} + +// The kernel has an attribute. +void calls_kernel_4() { + // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_5vE7Kernel5 + // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic + sycl::kernel_single_task([]() [[intel::named_sub_group_size(automatic)]] { // #Kernel5 + // expected-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}} + // expected-note@#Kernel5{{conflicting attribute is here}} + AttrFunc(); + }); +} +#endif // SYCL2020 diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp new file mode 100644 index 0000000000000..fa1fec97137d0 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp @@ -0,0 +1,55 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat -sycl-std=2017 -DSYCL2017 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat -sycl-std=2020 -DSYCL2020 %s + +// Tests for AST of sycl_explicit_simd function attribute. + +#include "sycl.hpp" + +sycl::queue deviceQueue; + +struct FuncObj { + [[intel::sycl_explicit_simd]] void operator()() const {} +}; + +[[intel::sycl_explicit_simd]] void func() {} + +int main() { + deviceQueue.submit([&](sycl::handler &h) { + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + h.single_task( + FuncObj()); + + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + h.single_task( + []() [[intel::sycl_explicit_simd]]{}); + +#if defined(SYCL2017) + // Test attribute is propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + // CHECK-NEXT: SYCLSimdAttr {{.*}} + h.single_task( + []() [[intel::sycl_explicit_simd]] { func(); }); +#endif //SYCL2017 + +#if defined(SYCL2020) + // Test attribute is not propagated. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + // CHECK-NEXT-NOT: SYCLSimdAttr {{.*}} + h.single_task( + []() [[intel::sycl_explicit_simd]] { func(); }); +#endif // SYCL2020 + }); + return 0; +} diff --git a/clang/test/SemaSYCL/sycl-esimd.cpp b/clang/test/SemaSYCL/sycl-esimd.cpp index 4c8d9da02e91d..5153aa1ea3b48 100644 --- a/clang/test/SemaSYCL/sycl-esimd.cpp +++ b/clang/test/SemaSYCL/sycl-esimd.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify %s // This test checks specifics of semantic analysis of ESIMD kernels. From a61aea03aa2e79f6ce329889e631353b6b24d02e Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Sat, 29 May 2021 07:09:43 -0700 Subject: [PATCH 2/6] update patch Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 17 +++++++---------- ...up-size-ast.cpp => named_sub_group_size.cpp} | 6 +----- 2 files changed, 8 insertions(+), 15 deletions(-) rename clang/test/SemaSYCL/{sub-group-size-ast.cpp => named_sub_group_size.cpp} (93%) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index c02a540b15c2f..5439b940ee6d1 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -328,14 +328,6 @@ void Sema::CheckDeprecatedSYCLAttributeSpelling(const ParsedAttr &A, return; } - // Diagnose the [[intel::named_sub_group_size]] attribute spelling - // in earlier SYCL mode (SYCL 2017) as being an ignored attribute. - if (LangOpts.getSYCLVersion() == LangOptions::SYCL_2017 && - A.getKind() == ParsedAttr::AT_IntelNamedSubGroupSize) { - Diag(A.getLoc(), diag::warn_attribute_ignored) << A; - return; - } - // Diagnose SYCL 2017 spellings in later SYCL modes. if (LangOpts.getSYCLVersion() > LangOptions::SYCL_2017) { // All attributes in the cl vendor namespace are deprecated in favor of a @@ -3362,8 +3354,6 @@ Sema::MergeIntelNamedSubGroupSizeAttr(Decl *D, static void handleIntelNamedSubGroupSize(Sema &S, Decl *D, const ParsedAttr &AL) { - S.CheckDeprecatedSYCLAttributeSpelling(AL); - StringRef SizeStr; SourceLocation Loc; if (AL.isArgIdent(0)) { @@ -3379,6 +3369,13 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D, SizeType)) { S.Diag(Loc, diag::warn_attribute_type_not_supported) << AL << SizeStr; } + + // If the [[intel::named_sub_group_size]] attribute spelling is used + // in earlier SYCL mode (SYCL 2017), we want to diagnose it as being + // an ignored attribute. + if (S.LangOpts.getSYCLVersion() == LangOptions::SYCL_2017) + S.Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL; + D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL)); } diff --git a/clang/test/SemaSYCL/sub-group-size-ast.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp similarity index 93% rename from clang/test/SemaSYCL/sub-group-size-ast.cpp rename to clang/test/SemaSYCL/named_sub_group_size.cpp index 5f3cb55e2bf2b..07de7e84f13d1 100644 --- a/clang/test/SemaSYCL/sub-group-size-ast.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -1,12 +1,8 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -DSYCL2017 -verify %s - // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,integer -DSYCL2020 %s // RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,primary -DSYCL2020 %s // Validate the semantic analysis checks for the named_sub_group_size attribute. -// The attribute can be only be applied to non-sycl-kernel/ -// non-sycl-device functions if they match the kernel they are being called -// from. #include "Inputs/sycl.hpp" @@ -49,7 +45,7 @@ void calls_kernel_2() { sycl::kernel_single_task(F1); } -// Test ttribute does not get propgated to the kernel. +// Test ttribute does not get propgated to the kernel. [[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc void calls_kernel_3() { From c887413de70e7ed4b1a6befa3740f481c384a5df Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 1 Jun 2021 20:42:23 -0700 Subject: [PATCH 3/6] address review comments Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 65 ++--- clang/lib/Sema/SemaDeclAttr.cpp | 4 +- clang/lib/Sema/SemaSYCL.cpp | 7 +- .../CodeGenSYCL/disable_loop_pipelining.cpp | 1 + .../test/CodeGenSYCL/initiation_interval.cpp | 1 + clang/test/CodeGenSYCL/intel-restrict.cpp | 2 +- .../test/CodeGenSYCL/loop_fuse_ind_device.cpp | 2 +- clang/test/CodeGenSYCL/max-concurrency.cpp | 1 + .../check-direct-attribute-propagation.cpp | 241 ++++++++++++++++++ .../test/SemaSYCL/disable_loop_pipelining.cpp | 2 + .../test/SemaSYCL/initiation_interval_ast.cpp | 1 + .../intel-fpga-no-global-work-offset.cpp | 24 +- .../intel-max-global-work-dim-device.cpp | 16 +- .../intel-max-work-group-size-device.cpp | 21 +- clang/test/SemaSYCL/intel-restrict.cpp | 18 +- clang/test/SemaSYCL/max-concurrency.cpp | 3 +- clang/test/SemaSYCL/max_global_work_dim.cpp | 97 +++++++ .../SemaSYCL/named_sub_group_size-ignore.cpp | 6 + clang/test/SemaSYCL/named_sub_group_size.cpp | 22 +- .../SemaSYCL/parallel_for_wrapper_attr.cpp | 2 +- .../SemaSYCL/reqd-sub-group-size-device.cpp | 41 +-- .../scheduler_target_fmax_mhz_ast.cpp | 17 +- clang/test/SemaSYCL/sycl-esimd-ast.cpp | 16 +- 23 files changed, 408 insertions(+), 202 deletions(-) create mode 100644 clang/test/SemaSYCL/check-direct-attribute-propagation.cpp create mode 100644 clang/test/SemaSYCL/max_global_work_dim.cpp create mode 100644 clang/test/SemaSYCL/named_sub_group_size-ignore.cpp diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index f8e3608720943..bf422af664150 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -416,11 +416,9 @@ def SYCLSimdDocs : Documentation { and code generation pipeline. Also, this attribute is used to distinguish ESIMD private globals from regular SYCL global variables. - When conforming to SYCL 1.2.1 extension specifications, when called from a - device kernel, the ``intel::sycl_explicit_simd`` is propagated from the - function it is applied to, to the kernel. When conforming to SYCL 2020 - extension specifications, the attribute is not propagated to the kernel. - + In SYCL 1.2.1 mode, the ``intel::sycl_explicit_simd`` attribute is propagated + from the function it is applied to onto the kernel which calls the function. + In SYCL 2020 mode, the attribute is not propagated to the kernel. }]; } @@ -2449,10 +2447,9 @@ lambda capture, or function object member, of the callable to which the attribute was applied. This effect is equivalent to annotating restrict on **all** kernel pointer arguments in an OpenCL or SPIR-V kernel. -When conforming to SYCL 1.2.1 extension specifications, when called from a -device kernel, the ``intel::kernel_args_restrict`` is propagated from the -function it is applied to, to the kernel. When conforming to SYCL 2020 -extension specifications, the attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::kernel_args_restrict`` attribute is propagated +from the function it is applied to onto the kernel which calls the function. +In SYCL 2020 mode, the attribute is not propagated to the kernel. The attribute forms an unchecked assertion, in that implementations do not need to check/confirm the pre-condition in any way. If a user applies @@ -2491,10 +2488,9 @@ def SYCLIntelNumSimdWorkItemsAttrDocs : Documentation { Applies to a device function/lambda function. Indicates the number of work items that should be processed in parallel. Valid values are positive integers. -When conforming to SYCL 1.2.1 extension specifications, when called from a -device kernel, the ``intel::num_simd_work_items`` is propagated from the -function it is applied to, to the kernel. When conforming to SYCL 2020 -extension specifications, the attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::num_simd_work_items`` attribute is propagated +from the function it is applied to onto the kernel which calls the function. +In SYCL 2020 mode, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -2644,10 +2640,10 @@ allows the Y and Z arguments to be optional. If not provided by the user, the value of Y and Z defaults to 1. See section 5.8.1 Kernel Attributes for more details. -When conforming to SYCL 1.2.1 specifications, when called from a device kernel, -the ``intel::reqd_work_group_size`` or [[sycl::reqd_work_group_size() is -propagated from the function it is applied to, to the kernel. When conforming -to SYCL 2020 specifications, the attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size`` or +``sycl::reqd_work_group_size`` attribute is propagated from the function it is +applied to onto the kernel which calls the function. In SYCL 2020 mode, the +attribute is not propagated to the kernel. .. code-block:: c++ @@ -2790,10 +2786,9 @@ of a work group. Values must be positive integers. This is similar to reqd_work_group_size, but allows work groups that are smaller or equal to the specified sizes. -When conforming to SYCL 1.2.1 specifications, when called from a device kernel, -the ``intel::max_work_group_size`` is propagated from the function it is -applied to, to the kernel. When conforming to SYCL 2020 specifications, the -attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::max_work_group_size`` attribute is propagated +from the function it is applied to onto the kernel which calls the function. +In SYCL 2020 mode, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -2825,10 +2820,9 @@ function object). Indicates the largest valid global work dimension that will be accepted when running the kernel on a device. Valid values are integers in a range of [0, 3]. -When conforming to SYCL 1.2.1 extension specifications, when called from a -device kernel, the ``intel::max_global_work_dim`` is propagated from the -function it is applied to, to the kernel. When conforming to SYCL 2020 -extension specifications, the attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::max_global_work_dim`` attribute is propagated +from the function it is applied to onto the kernel which calls the function. +In SYCL 2020 mode, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -2885,10 +2879,9 @@ device operation, guiding the FPGA backend to insert the appropriate number of registers to break-up the combinational logic circuit, and thereby controlling the length of the longest combinational path. -When conforming to SYCL 1.2.1 extension specifications, when called from a -device kernel, the ``intel::scheduler_target_fmax_mhz`` is propagated from the -function it is applied to, to the kernel. When conforming to SYCL 2020 -extension specifications, the attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::scheduler_target_fmax_mhz`` attribute is +propagated from the function it is applied to onto the kernel which calls the +function. In SYCL 2020 mode, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -2920,10 +2913,9 @@ function object). If 1, compiler doesn't use the global work offset values for the device function. Valid values are 0 and 1. If used without argument, value of 1 is set implicitly. -When conforming to SYCL 1.2.1 extension specifications, when called from a -device kernel, the ``intel::no_global_work_offset`` is propagated from the -function it is applied to, to the kernel. When conforming to SYCL 2020 -extension specifications, the attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::no_global_work_offset`` attribute is +propagated from the function it is applied to onto the kernel which calls the +function. In SYCL 2020 mode, the attribute is not propagated to the kernel. .. code-block:: c++ @@ -4599,10 +4591,9 @@ In addition to device functions, the required sub-group size attribute may also be specified in the definition of a named functor object and lambda functions, as in the examples below: -When conforming to SYCL 1.2.1 specifications, when called from a device kernel, -the ``intel::reqd_sub_group_size`` is propagated from the function it is -applied to, to the kernel. When conforming to SYCL 2020 specifications, the -attribute is not propagated to the kernel. +In SYCL 1.2.1 mode, the ``intel::reqd_sub_group_size`` attribute is propagated +from the function it is applied to onto the kernel which calls the function. +In SYCL 2020 mode, the attribute is not propagated to the kernel. .. code-block:: c++ diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 5439b940ee6d1..f07d2d26ba7b3 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3373,8 +3373,10 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D, // If the [[intel::named_sub_group_size]] attribute spelling is used // in earlier SYCL mode (SYCL 2017), we want to diagnose it as being // an ignored attribute. - if (S.LangOpts.getSYCLVersion() == LangOptions::SYCL_2017) + if (S.LangOpts.getSYCLVersion() == LangOptions::SYCL_2017) { S.Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL; + return; + } D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL)); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 76ccb30e09c46..f40deeef7b96e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -358,8 +358,11 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, // Attributes that should not be propagated from device functions to a // kernel in SYCL 1.2.1. if (DirectlyCalled) { - llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), - [](Attr *A) { return isa(A); }); + llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { + return isa(A); + }); } } else { // Attributes that should not be propagated from device functions to a diff --git a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp index 97c4c754896b5..2912466b4fefd 100644 --- a/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp +++ b/clang/test/CodeGenSYCL/disable_loop_pipelining.cpp @@ -1,3 +1,4 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/initiation_interval.cpp b/clang/test/CodeGenSYCL/initiation_interval.cpp index 0fa5699b32e99..df7f0f5ad6299 100644 --- a/clang/test/CodeGenSYCL/initiation_interval.cpp +++ b/clang/test/CodeGenSYCL/initiation_interval.cpp @@ -1,3 +1,4 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/intel-restrict.cpp b/clang/test/CodeGenSYCL/intel-restrict.cpp index 20aa089e35f26..343616ccc2d54 100644 --- a/clang/test/CodeGenSYCL/intel-restrict.cpp +++ b/clang/test/CodeGenSYCL/intel-restrict.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -triple spir64-unknown-unknown-sycldevice -o - | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device %s -emit-llvm -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -o - | FileCheck %s template __attribute__((sycl_kernel)) void kernel(const Func &kernelFunc) { diff --git a/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp b/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp index ca3eb4445e1bd..b8c6c3c8e7254 100644 --- a/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp +++ b/clang/test/CodeGenSYCL/loop_fuse_ind_device.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/max-concurrency.cpp b/clang/test/CodeGenSYCL/max-concurrency.cpp index 5bafcd12ab89e..bd0bc5116c7d4 100644 --- a/clang/test/CodeGenSYCL/max-concurrency.cpp +++ b/clang/test/CodeGenSYCL/max-concurrency.cpp @@ -1,3 +1,4 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -disable-llvm-passes -triple spir64-unknown-unknown-sycldevice -sycl-std=2020 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp new file mode 100644 index 0000000000000..9f86c5521fa92 --- /dev/null +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -0,0 +1,241 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -ast-dump %s | FileCheck %s + +// Tests to validate the SYCL 2020 requirement mandating the avoidance of the propagation of all kernel attributes to the caller when used on a function. + +#include "sycl.hpp" + +sycl::queue deviceQueue; + +struct FuncObj { + [[intel::sycl_explicit_simd]] void operator()() const {} +}; + +struct FuncObj1 { + [[intel::no_global_work_offset(1)]] void operator()() const {} +}; + +struct FuncObj2 { + [[intel::scheduler_target_fmax_mhz(10)]] void operator()() const {} +}; + +struct FuncObj3 { + [[intel::max_work_group_size(2, 2, 2)]] void operator()() const {} +}; + +struct FuncObj4 { + [[intel::reqd_work_group_size(2, 2, 2)]] void operator()() const {} +}; + +struct FuncObj5 { + [[intel::num_simd_work_items(8)]] void operator()() const {} +}; + +struct FuncObj6 { + [[intel::kernel_args_restrict]] void operator()() const {} +}; + +[[intel::sycl_explicit_simd]] void func() {} + +[[intel::no_global_work_offset(1)]] void func1() {} + +[[intel::scheduler_target_fmax_mhz(2)]] void func2() {} + +[[intel::max_work_group_size(1, 1, 1)]] void func3() {} + +[[intel::reqd_work_group_size(1, 1, 1)]] void func4() {} + +[[intel::num_simd_work_items(5)]] void func5() {} + +[[intel::kernel_args_restrict]] void func6() {} + +int main() { + deviceQueue.submit([&](sycl::handler &h) { + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + h.single_task( + FuncObj()); + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + h.single_task( + []() [[intel::sycl_explicit_simd]]{}); + + // Test attribute is not propagated from function. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + // CHECK-NOT: SYCLSimdAttr {{.*}} + h.single_task( + []() [[intel::sycl_explicit_simd]] { func(); }); + + // Test attribute is not propagated from function. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 'void ()' + // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + h.single_task( + []() { func1(); }); + + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5 'void ()' + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task( + FuncObj1()); + + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kerne6 'void ()' + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task( + []() [[intel::no_global_work_offset]]{}); + + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel7 'void () + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 10 + // CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} + h.single_task( + FuncObj2()); + + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel8 'void ()' + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 20 + // CHECK-NEXT: IntegerLiteral{{.*}}20{{$}} + h.single_task( + []() [[intel::scheduler_target_fmax_mhz(20)]]{}); + + // Test attribute is not propagated from function. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9 'void ()' + // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + h.single_task( + []() { func2(); }); + + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10 'void ()' + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + h.single_task( + FuncObj3()); + + // Test attribute is not propagated from function. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 'void ()' + // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + h.single_task( + []() { func3(); }); + + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 'void ()' + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + h.single_task( + []() [[intel::max_work_group_size(8, 8, 8)]]{}); + + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel13 'void ()' + // CHECK: ReqdWorkGroupSizeAttr{{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + h.single_task( + FuncObj4()); + + // Test attribute is not propagated from function. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel14 'void ()' + // CHECK-NOT: ReqdWorkGroupSizeAttr {{.*}} + h.single_task( + []() { func4(); }); + + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel15 'void ()' + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + h.single_task( + []() [[intel::reqd_work_group_size(8, 8, 8)]]{}); + + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel16 'void () + // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + h.single_task( + FuncObj5()); + + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel17 'void ()' + // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 20 + // CHECK-NEXT: IntegerLiteral{{.*}}20{{$}} + h.single_task( + []() [[intel::num_simd_work_items(20)]]{}); + + // Test attribute is not propagated from function. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel18 'void ()' + // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} + h.single_task( + []() { func5(); }); + + // Test attribute directly applies on kernel functor. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel19 'void () + // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} + h.single_task( + FuncObj6()); + + // Test attribute directly applies on kernel lambda. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel20 'void ()' + // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} + h.single_task( + []() [[intel::kernel_args_restrict]]{}); + + // Test attribute is not propagated from functiom. + // CHECK-LABEL: FunctionDecl {{.*}}test_kernel21 'void ()' + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr {{.*}} + h.single_task( + []() { func6(); }); + + }); + return 0; +} diff --git a/clang/test/SemaSYCL/disable_loop_pipelining.cpp b/clang/test/SemaSYCL/disable_loop_pipelining.cpp index c36f0f6d29cdb..c4b37d168c97f 100644 --- a/clang/test/SemaSYCL/disable_loop_pipelining.cpp +++ b/clang/test/SemaSYCL/disable_loop_pipelining.cpp @@ -1,3 +1,5 @@ +// RUN: %clang_cc1 %s -fsyntax-only -internal-isystem %S/Inputs -fsycl-is-device -Wno-sycl-2017-compat -sycl-std=2017 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat -sycl-std=2017 %s | FileCheck %s // RUN: %clang_cc1 %s -fsyntax-only -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -DTRIGGER_ERROR -verify // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -sycl-std=2020 %s | FileCheck %s diff --git a/clang/test/SemaSYCL/initiation_interval_ast.cpp b/clang/test/SemaSYCL/initiation_interval_ast.cpp index 777230e224572..cfcbc131299a7 100644 --- a/clang/test/SemaSYCL/initiation_interval_ast.cpp +++ b/clang/test/SemaSYCL/initiation_interval_ast.cpp @@ -1,3 +1,4 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s // Tests for AST of Intel FPGA initiation_interval function attributes. diff --git a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index e342d861385ee..5851c27fac328 100644 --- a/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp +++ b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp @@ -1,5 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -sycl-std=2017 -ast-dump -DSYCL2017 -verify -pedantic %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -sycl-std=2020 -ast-dump -DSYCL2020 -verify -pedantic %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-return-type -sycl-std=2017 -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s #include "sycl.hpp" @@ -12,8 +11,6 @@ struct FuncObj { [[intelfpga::no_global_work_offset]] void operator()() const {} }; -[[intel::no_global_work_offset(1)]] void func() {} - int main() { q.submit([&](handler &h) { // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} @@ -65,25 +62,6 @@ int main() { h.single_task( []() [[intel::no_global_work_offset(0), // expected-note {{previous attribute is here}} intel::no_global_work_offset(1)]]{}); // expected-warning{{attribute 'no_global_work_offset' is already applied with different arguments}} - -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9 - // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - h.single_task( - []() { func(); }); -#endif // SYCL2020 - -#if defined(SYCL2017) - // Test attribute is propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kerne10 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - h.single_task( - []() { func(); }); -#endif // SYCL2017 }); return 0; } diff --git a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp index 2c37f0ec52238..7b1171147b113 100644 --- a/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp +++ b/clang/test/SemaSYCL/intel-max-global-work-dim-device.cpp @@ -1,7 +1,5 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s #include "sycl.hpp" @@ -251,16 +249,7 @@ int main() { // expected-note@+2 {{did you mean to use 'intel::max_global_work_dim' instead?}} h.single_task( []() [[intelfpga::max_global_work_dim(2)]]{}); -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr {{.*}} - h.single_task( - []() { func_do_not_ignore(); }); -#endif // SYCL2020 -#if defined(SYCL2017) - // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -475,7 +464,6 @@ int main() { // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 0 // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} -#endif // SYCL2017 // Ignore duplicate attribute with same argument value. h.single_task( diff --git a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp index ea7e2718ab0c0..02c7d487666dc 100644 --- a/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-max-work-group-size-device.cpp @@ -1,7 +1,5 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -Wno-sycl-2017-compat -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s #include "sycl.hpp" @@ -36,7 +34,6 @@ struct Func { }; #ifdef TRIGGER_ERROR -#if defined(SYCL2020) struct DAFuncObj { [[intel::max_work_group_size(4, 4, 4)]] [[cl::reqd_work_group_size(8, 8, 4)]] // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} \ @@ -44,7 +41,6 @@ struct DAFuncObj { // expected-note{{did you mean to use 'sycl::reqd_work_group_size' instead?}} void operator()() const {} }; -#endif // SYCL2020 #endif // TRIGGER_ERROR int main() { @@ -78,16 +74,6 @@ int main() { h.single_task( []() [[intelfpga::max_work_group_size(8, 8, 8)]]{}); -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} - h.single_task( - []() { func_do_not_ignore(); }); -#endif // SYCL2020 - -#if defined(SYCL2017) - // Test attribute is propagated // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -101,7 +87,6 @@ int main() { // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} h.single_task( []() { func_do_not_ignore(); }); -#endif // SYCL2017 // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} @@ -148,10 +133,8 @@ int main() { []() [[intel::max_work_group_size(16, 16, 16), // expected-note{{conflicting attribute is here}} intel::max_work_group_size(2, 2, 2)]]{}); // expected-warning{{attribute 'max_work_group_size' is already applied with different arguments}} -#if defined(SYCL2020) h.single_task( DAFuncObj()); -#endif // SYCL2020 #endif // TRIGGER_ERROR }); diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index 2970d601f2c02..f8e7a670fddd6 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -1,8 +1,5 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DCHECKDIAG -verify -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DCHECKDIAG -verify -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -triple spir64 -DSYCL2017 %s -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -triple spir64 -DSYCL2020 %s - +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 -DCHECKDIAG -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64 | FileCheck %s [[intel::kernel_args_restrict]] void func_do_not_ignore() {} @@ -29,19 +26,8 @@ int main() { kernel( []() [[intel::kernel_args_restrict]] {}); -#if defined(SYCL2017) - // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLIntelKernelArgsRestrictAttr kernel( []() { func_do_not_ignore(); }); -#endif // SYCL2017 - -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 - // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr - kernel( - []() { func_do_not_ignore(); }); -#endif // SYCL2020 } diff --git a/clang/test/SemaSYCL/max-concurrency.cpp b/clang/test/SemaSYCL/max-concurrency.cpp index a10045413711b..b4382921aa70d 100644 --- a/clang/test/SemaSYCL/max-concurrency.cpp +++ b/clang/test/SemaSYCL/max-concurrency.cpp @@ -1,4 +1,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/max_global_work_dim.cpp b/clang/test/SemaSYCL/max_global_work_dim.cpp new file mode 100644 index 0000000000000..8d921498d91f5 --- /dev/null +++ b/clang/test/SemaSYCL/max_global_work_dim.cpp @@ -0,0 +1,97 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s + +// Tests for AST of Intel FPGA max_global_work_dim function attribute in SYCL 2020 mode. +#include "sycl.hpp" + +sycl::queue deviceQueue; + +// CHECK: FunctionDecl {{.*}} func1 'void ()' +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: SYCLIntelMaxGlobalWorkDimAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 2 +// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 +[[intel::max_global_work_dim(2)]] void func1() {} + +// Test that checks template parameter support on function. +// CHECK: FunctionTemplateDecl {{.*}} func2 +// CHECK: FunctionDecl {{.*}} func2 'void ()' +// CHECK-NEXT: CompoundStmt +// CHECK_NEXT: SYCLIntelMaxGlobalWorkDimAttr {{.*}} +// CHECK_NEXT: DeclRefExpr {{.*}} 'int' NonTypeTemplateParm {{.*}} 'N' 'int' +// CHECK: FunctionDecl {{.*}} func2 'void ()' +// CHECK-NEXT: TemplateArgument integral 2 +// CHECK-NEXT: CompoundStmt +// CHECK-NEXT: SYCLIntelMaxGlobalWorkDimAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 2 +// CHECK-NEXT: SubstNonTypeTemplateParmExpr +// CHECK-NEXT: NonTypeTemplateParmDecl +// CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 +template +[[intel::max_global_work_dim(N)]] void func2() {} + +// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored. +// CHECK: FunctionDecl {{.*}} {{.*}} func3 'void ()' +// CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} +// CHECK-NEXT: ConstantExpr {{.*}} 'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +[[intel::max_global_work_dim(1)]] +[[intel::max_global_work_dim(1)]] void func3() {} + +class KernelFunctor { +public: + void operator()() const { + func1(); + } +}; + +// Test that checks template parameter support on class member function. +template +class KernelFunctor2 { +public: + [[intel::max_global_work_dim(N)]] void operator()() const { + } +}; + +int main() { + deviceQueue.submit([&](sycl::handler &h) { + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 + // CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr + KernelFunctor f1; + h.single_task(f1); + + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2 + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 2 + // CHECK-NEXT: SubstNonTypeTemplateParmExpr + // CHECK-NEXT: NonTypeTemplateParmDecl + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 + KernelFunctor2<2> f2; + h.single_task(f2); + + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_3 + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 + h.single_task( + []() [[intel::max_global_work_dim(1)]]{}); + + // Ignore duplicate attribute. + h.single_task( + // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4 + // CHECK: SYCLIntelMaxGlobalWorkDimAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + []() [[intel::max_global_work_dim(1), + intel::max_global_work_dim(1)]]{}); + }); + + func2<2>(); + + return 0; +} diff --git a/clang/test/SemaSYCL/named_sub_group_size-ignore.cpp b/clang/test/SemaSYCL/named_sub_group_size-ignore.cpp new file mode 100644 index 0000000000000..9cf0d0634a8b2 --- /dev/null +++ b/clang/test/SemaSYCL/named_sub_group_size-ignore.cpp @@ -0,0 +1,6 @@ +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -verify %s + +// Test that we get ignored attribute warning when using +// a [[intel::named_sub_group_size()]] attribute spelling while not +// in SYCL 2020 mode. +[[intel::named_sub_group_size(automatic)]] void func_ignore(); // expected-warning {{'named_sub_group_size' attribute ignored}} diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index 07de7e84f13d1..05ad83b15cab6 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -1,19 +1,10 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -DSYCL2017 -verify %s -// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,integer -DSYCL2020 %s -// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,primary -DSYCL2020 %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=primary -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,integer %s | FileCheck %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -fsycl-default-sub-group-size=10 -sycl-std=2020 -internal-isystem %S/Inputs -fsyntax-only -ast-dump -verify=expected,primary %s | FileCheck %s -// Validate the semantic analysis checks for the named_sub_group_size attribute. +// Validate the semantic analysis checks for the named_sub_group_size attribute in SYCL 2020 mode. #include "Inputs/sycl.hpp" -#if defined(SYCL2017) -// Test that we get ignored attribute warning when using -// a [[intel::named_sub_group_size()]] attribute spelling while not -// in SYCL 2020 mode. -[[intel::named_sub_group_size(automatic)]] void func_ignore(); // expected-warning {{'named_sub_group_size' attribute ignored}} -#endif // SYCL2017 - -#if defined(SYCL2020) // The kernel has an attribute. void calls_kernel_1() { // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_1vE7Kernel1 @@ -40,7 +31,7 @@ void calls_kernel_2() { sycl::kernel_single_task(F); Functor1 F1; - // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_3vE7Kernel3 + // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_2vE7Kernel3 'void ()' // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Primary sycl::kernel_single_task(F1); } @@ -49,7 +40,7 @@ void calls_kernel_2() { [[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc void calls_kernel_3() { - // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_4vE7Kernel4 + // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_3vE7Kernel4 'void ()' // CHECK-NOT : IntelNamedSubGroupSizeAttr {{.*}} sycl::kernel_single_task([]() { // #Kernel4 // primary-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}} @@ -60,7 +51,7 @@ void calls_kernel_3() { // The kernel has an attribute. void calls_kernel_4() { - // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_5vE7Kernel5 + // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_4vE7Kernel5 'void ()' // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic sycl::kernel_single_task([]() [[intel::named_sub_group_size(automatic)]] { // #Kernel5 // expected-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}} @@ -68,4 +59,3 @@ void calls_kernel_4() { AttrFunc(); }); } -#endif // SYCL2020 diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp index dfc39bc84a6d2..ddf80076e14e4 100755 --- a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp +++ b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s +// RUN: %clang_cc1 %s -fsyntax-only -sycl-std=2017 -ast-dump -fsycl-is-device -triple spir64 | FileCheck %s #include "Inputs/sycl.hpp" diff --git a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp index 36a269a456f1d..d139392620551 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp @@ -1,25 +1,14 @@ -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -DTRIGGER_ERROR -DSYCL2017 -verify %s -// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -DTRIGGER_ERROR -DSYCL2020 -verify %s -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -DSYCL2017 %s -// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -DSYCL2020 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s #include "sycl.hpp" using namespace cl::sycl; queue q; -[[intel::reqd_sub_group_size(16)]] void func() {} - -#if defined(SYCL2017) [[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} [[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}} -#endif // SYCL2017 - -#if defined(SYCL2020) -[[intel::reqd_sub_group_size(4)]] void foo() {} // OK -[[intel::reqd_sub_group_size(32)]] void baz() {} // OK -#endif // SYCL2020 // No diagnostic is emitted because the arguments match. [[intel::reqd_sub_group_size(12)]] void bar(); @@ -43,7 +32,6 @@ class Functor16 { [[intel::reqd_sub_group_size(16)]] void operator()() const {} }; -#if defined(SYCL2017) class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} public: [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} @@ -51,13 +39,11 @@ class Functor8 { // expected-error {{conflicting attributes applied to a SYCL ke } }; -#endif // SYCL2017 class Functor4 { public: [[intel::reqd_sub_group_size(12)]] void operator()() const {} }; -#if defined(SYCL2017) class Functor { public: void operator()() const { @@ -65,15 +51,11 @@ class Functor { } }; -#endif // SYCL2017 - int main() { q.submit([&](handler &h) { Functor16 f16; h.single_task(f16); -#if defined(SYCL2017) - // Test attribute is propagated. Functor f; h.single_task(f); @@ -86,7 +68,6 @@ int main() { baz(); }); #endif // TRIGGER_ERROR -#endif // SYCL2017 h.single_task([]() [[intel::reqd_sub_group_size(2)]]{}); h.single_task([]() [[intel::reqd_sub_group_size(4)]] { foo(); }); @@ -94,24 +75,6 @@ int main() { Functor4 f4; h.single_task(f4); - -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}class kernel_name9 - // CHECK-NOT: IntelReqdSubGroupSizeAttr {{.*}} - h.single_task( - []() { func(); }); -#endif // SYCL2020 - -#if defined(SYCL2017) - // Test attribute is propagated. - // CHECK-LABEL: FunctionDecl {{.*}}class kernel_name10 - // CHECK: IntelReqdSubGroupSizeAttr {{.*}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 16 - h.single_task( - []() { func(); }); -#endif // SYCL2017 }); return 0; } diff --git a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp index e3f53b23f5807..19f22d391a72a 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp @@ -1,5 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2017 -ast-dump -DSYCL2017 %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -ast-dump -DSYCL2020 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s // Tests for AST of Intel FPGA scheduler_target_fmax_mhz function attribute. #include "sycl.hpp" @@ -62,23 +61,12 @@ class KernelFunctor2 { int main() { deviceQueue.submit([&](sycl::handler &h) { -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 - // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr - KernelFunctor f1; - h.single_task(f1); -#endif // SYCL2020 - -#if defined(SYCL2017) // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr KernelFunctor f1; h.single_task(f1); -#endif // SYCL2017 - // Test class template argument. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_2 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' @@ -97,8 +85,6 @@ int main() { h.single_task( []() [[intel::scheduler_target_fmax_mhz(4)]]{}); -#if defined(SYCL2017) - // Test function template argument. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_4 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' @@ -108,7 +94,6 @@ int main() { // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 75 h.single_task( []() { func3<75>(); }); -#endif // SYCL2017 // Ignore duplicate attribute. h.single_task( diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp index fa1fec97137d0..5343019f43a08 100644 --- a/clang/test/SemaSYCL/sycl-esimd-ast.cpp +++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp @@ -1,5 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat -sycl-std=2017 -DSYCL2017 %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -ast-dump -Wno-sycl-2017-compat -sycl-std=2020 -DSYCL2020 %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s // Tests for AST of sycl_explicit_simd function attribute. @@ -29,7 +28,6 @@ int main() { h.single_task( []() [[intel::sycl_explicit_simd]]{}); -#if defined(SYCL2017) // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLSimdAttr {{.*}} Implicit @@ -38,18 +36,6 @@ int main() { // CHECK-NEXT: SYCLSimdAttr {{.*}} h.single_task( []() [[intel::sycl_explicit_simd]] { func(); }); -#endif //SYCL2017 - -#if defined(SYCL2020) - // Test attribute is not propagated. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr {{.*}} - // CHECK-NEXT-NOT: SYCLSimdAttr {{.*}} - h.single_task( - []() [[intel::sycl_explicit_simd]] { func(); }); -#endif // SYCL2020 }); return 0; } From 712c32bb0c3830f5be64998df109639ee0498ef0 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 1 Jun 2021 20:46:46 -0700 Subject: [PATCH 4/6] address review comments Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaDeclAttr.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index f07d2d26ba7b3..a866953fe8984 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3370,9 +3370,8 @@ static void handleIntelNamedSubGroupSize(Sema &S, Decl *D, S.Diag(Loc, diag::warn_attribute_type_not_supported) << AL << SizeStr; } - // If the [[intel::named_sub_group_size]] attribute spelling is used - // in earlier SYCL mode (SYCL 2017), we want to diagnose it as being - // an ignored attribute. + // If the [[intel::named_sub_group_size]] attribute spelling is used in + // SYCL 2017 mode, we want to diagnose it as being an ignored attribute. if (S.LangOpts.getSYCLVersion() == LangOptions::SYCL_2017) { S.Diag(AL.getLoc(), diag::warn_attribute_ignored) << AL; return; From a62d6ccd15dbe5787b9da446f4916f395d1df586 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 1 Jun 2021 20:50:22 -0700 Subject: [PATCH 5/6] fix clang format issues Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaSYCL.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index f40deeef7b96e..a0aed01c7dd63 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -360,7 +360,7 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, if (DirectlyCalled) { llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { return isa(A); }); } From ade77f5f69666a2042c88abce6a0e4391f5922dc Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 2 Jun 2021 06:15:18 -0700 Subject: [PATCH 6/6] fix test failures Signed-off-by: Soumi Manna --- .../check-direct-attribute-propagation.cpp | 74 +++++++++---------- clang/test/SemaSYCL/named_sub_group_size.cpp | 20 ++--- 2 files changed, 47 insertions(+), 47 deletions(-) diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp index 9f86c5521fa92..80fcf875afc6a 100644 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -51,14 +51,14 @@ struct FuncObj6 { int main() { deviceQueue.submit([&](sycl::handler &h) { // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel1 + // CHECK: FunctionDecl {{.*}}test_kernel1 // CHECK: SYCLSimdAttr {{.*}} Implicit // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit // CHECK-NEXT: SYCLSimdAttr {{.*}} h.single_task( FuncObj()); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel2 + // CHECK: FunctionDecl {{.*}}test_kernel2 // CHECK: SYCLSimdAttr {{.*}} Implicit // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit // CHECK-NEXT: SYCLSimdAttr {{.*}} @@ -66,31 +66,31 @@ int main() { []() [[intel::sycl_explicit_simd]]{}); // Test attribute is not propagated from function. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr {{.*}} - // CHECK-NOT: SYCLSimdAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel3 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr {{.*}} + // CHECK-NOT: SYCLSimdAttr {{.*}} h.single_task( []() [[intel::sycl_explicit_simd]] { func(); }); // Test attribute is not propagated from function. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel4 'void ()' - // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel4 + // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} h.single_task( []() { func1(); }); // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel5 'void ()' - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: FunctionDecl {{.*}}test_kernel5 + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} h.single_task( FuncObj1()); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kerne6 'void ()' + // CHECK: FunctionDecl {{.*}}test_kerne6 // CHECK: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 @@ -99,7 +99,7 @@ int main() { []() [[intel::no_global_work_offset]]{}); // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel7 'void () + // CHECK: FunctionDecl {{.*}}test_kernel7 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 10 @@ -108,7 +108,7 @@ int main() { FuncObj2()); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel8 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel8 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 20 @@ -117,13 +117,13 @@ int main() { []() [[intel::scheduler_target_fmax_mhz(20)]]{}); // Test attribute is not propagated from function. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel9 'void ()' - // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel9 + // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} h.single_task( []() { func2(); }); // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel10 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel10 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 @@ -138,13 +138,13 @@ int main() { FuncObj3()); // Test attribute is not propagated from function. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel11 'void ()' - // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel11 + // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} h.single_task( []() { func3(); }); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel12 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel12 // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 @@ -159,7 +159,7 @@ int main() { []() [[intel::max_work_group_size(8, 8, 8)]]{}); // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel13 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel13 // CHECK: ReqdWorkGroupSizeAttr{{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 @@ -174,13 +174,13 @@ int main() { FuncObj4()); // Test attribute is not propagated from function. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel14 'void ()' - // CHECK-NOT: ReqdWorkGroupSizeAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel14 + // CHECK-NOT: ReqdWorkGroupSizeAttr {{.*}} h.single_task( []() { func4(); }); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel15 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel15 // CHECK: ReqdWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 @@ -195,7 +195,7 @@ int main() { []() [[intel::reqd_work_group_size(8, 8, 8)]]{}); // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel16 'void () + // CHECK: FunctionDecl {{.*}}test_kernel16 // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 8 @@ -204,7 +204,7 @@ int main() { FuncObj5()); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel17 'void ()' + // CHECK: FunctionDecl {{.*}}test_kernel17 // CHECK: SYCLIntelNumSimdWorkItemsAttr {{.*}} // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 20 @@ -213,26 +213,26 @@ int main() { []() [[intel::num_simd_work_items(20)]]{}); // Test attribute is not propagated from function. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel18 'void ()' - // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel18 + // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} h.single_task( []() { func5(); }); // Test attribute directly applies on kernel functor. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel19 'void () - // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel19 + // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} h.single_task( FuncObj6()); // Test attribute directly applies on kernel lambda. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel20 'void ()' - // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel20 + // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} h.single_task( []() [[intel::kernel_args_restrict]]{}); // Test attribute is not propagated from functiom. - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel21 'void ()' - // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr {{.*}} + // CHECK: FunctionDecl {{.*}}test_kernel21 + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr {{.*}} h.single_task( []() { func6(); }); diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index 05ad83b15cab6..13ce54c4f0116 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -7,8 +7,8 @@ // The kernel has an attribute. void calls_kernel_1() { - // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_1vE7Kernel1 - // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic + // CHECK: FunctionDecl {{.*}}Kernel1 + // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic sycl::kernel_single_task([]() [[intel::named_sub_group_size(automatic)]] { }); } @@ -26,13 +26,13 @@ struct Functor1 { // Test attributes get propgated to the kernel. void calls_kernel_2() { Functor F; - // CHECK-LABEL: FunctionDecl {{.*}} _ZTSZ14calls_kernel_2vE7Kernel2 - // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic + // CHECK: FunctionDecl {{.*}}Kernel2 + // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic sycl::kernel_single_task(F); Functor1 F1; - // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_2vE7Kernel3 'void ()' - // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Primary + // CHECK: FunctionDecl {{.*}}Kernel3 + // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Primary sycl::kernel_single_task(F1); } @@ -40,8 +40,8 @@ void calls_kernel_2() { [[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc void calls_kernel_3() { - // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_3vE7Kernel4 'void ()' - // CHECK-NOT : IntelNamedSubGroupSizeAttr {{.*}} + // CHECK: FunctionDecl {{.*}}Kernel4 + // CHECK-NOT: IntelNamedSubGroupSizeAttr {{.*}} sycl::kernel_single_task([]() { // #Kernel4 // primary-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}} // primary-note@#Kernel4{{kernel declared here}} @@ -51,8 +51,8 @@ void calls_kernel_3() { // The kernel has an attribute. void calls_kernel_4() { - // CHECK-LABEL: FunctionDecl {{.*}}_ZTSZ14calls_kernel_4vE7Kernel5 'void ()' - // CHECK : IntelNamedSubGroupSizeAttr {{.*}} Automatic + // CHECK: FunctionDecl {{.*}}Kernel5 + // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic sycl::kernel_single_task([]() [[intel::named_sub_group_size(automatic)]] { // #Kernel5 // expected-error@#AttrFunc{{kernel-called function must have a sub group size that matches the size specified for the kernel}} // expected-note@#Kernel5{{conflicting attribute is here}}