diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index e5eb377cbfe07..bf422af664150 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -415,6 +415,10 @@ 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. + + 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. }]; } @@ -2443,8 +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. -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. +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 @@ -2482,8 +2487,10 @@ 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. + +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++ @@ -2633,6 +2640,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. +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++ [[sycl::reqd_work_group_size(4, 4, 4)]] void foo() {} @@ -2773,8 +2785,10 @@ 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. + +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++ @@ -2805,8 +2819,10 @@ 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. + +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++ @@ -2863,6 +2879,10 @@ 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. +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++ [[intel::scheduler_target_fmax_mhz(4)]] void foo() {} @@ -2893,6 +2913,10 @@ 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. +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++ [[intel::no_global_work_offset]] @@ -4567,6 +4591,10 @@ 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: +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++ class Functor diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index dd9cc2d949c8b..a866953fe8984 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -3369,6 +3369,14 @@ 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 + // 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; + } + D->addAttr(IntelNamedSubGroupSizeAttr::Create(S.Context, SizeType, AL)); } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index c2c11765a8bc2..a0aed01c7dd63 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -343,15 +343,43 @@ 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 +394,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/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-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/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_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..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 -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/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-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp new file mode 100644 index 0000000000000..80fcf875afc6a --- /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: 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: 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: 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: FunctionDecl {{.*}}test_kernel4 + // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr {{.*}} + h.single_task( + []() { func1(); }); + + // Test attribute directly applies on kernel functor. + // 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: FunctionDecl {{.*}}test_kerne6 + // 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: FunctionDecl {{.*}}test_kernel7 + // 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: FunctionDecl {{.*}}test_kernel8 + // 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: FunctionDecl {{.*}}test_kernel9 + // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr {{.*}} + h.single_task( + []() { func2(); }); + + // Test attribute directly applies on kernel functor. + // CHECK: FunctionDecl {{.*}}test_kernel10 + // 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: FunctionDecl {{.*}}test_kernel11 + // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr {{.*}} + h.single_task( + []() { func3(); }); + + // Test attribute directly applies on kernel lambda. + // CHECK: FunctionDecl {{.*}}test_kernel12 + // 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: FunctionDecl {{.*}}test_kernel13 + // 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: FunctionDecl {{.*}}test_kernel14 + // CHECK-NOT: ReqdWorkGroupSizeAttr {{.*}} + h.single_task( + []() { func4(); }); + + // Test attribute directly applies on kernel lambda. + // CHECK: FunctionDecl {{.*}}test_kernel15 + // 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: FunctionDecl {{.*}}test_kernel16 + // 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: FunctionDecl {{.*}}test_kernel17 + // 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: FunctionDecl {{.*}}test_kernel18 + // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr {{.*}} + h.single_task( + []() { func5(); }); + + // Test attribute directly applies on kernel functor. + // CHECK: FunctionDecl {{.*}}test_kernel19 + // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} + h.single_task( + FuncObj6()); + + // Test attribute directly applies on kernel lambda. + // CHECK: FunctionDecl {{.*}}test_kernel20 + // CHECK: SYCLIntelKernelArgsRestrictAttr {{.*}} + h.single_task( + []() [[intel::kernel_args_restrict]]{}); + + // Test attribute is not propagated from functiom. + // CHECK: FunctionDecl {{.*}}test_kernel21 + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr {{.*}} + h.single_task( + []() { func6(); }); + + }); + return 0; +} 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/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 b3c2d4b8e40c2..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,4 +1,4 @@ -// 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 -sycl-std=2017 -Wno-sycl-2017-compat -fcxx-exceptions -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s #include "sycl.hpp" @@ -61,7 +61,7 @@ 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}} }); 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 29cc7f084e2d4..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,5 +1,5 @@ -// 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 -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" 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..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,5 +1,5 @@ -// 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 -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" 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..f8e7a670fddd6 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -1,5 +1,5 @@ -// 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 -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() {} 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/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 new file mode 100644 index 0000000000000..13ce54c4f0116 --- /dev/null +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -0,0 +1,61 @@ +// 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 in SYCL 2020 mode. + +#include "Inputs/sycl.hpp" + +// The kernel has an attribute. +void calls_kernel_1() { + // CHECK: FunctionDecl {{.*}}Kernel1 + // 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: FunctionDecl {{.*}}Kernel2 + // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic + sycl::kernel_single_task(F); + + Functor1 F1; + // CHECK: FunctionDecl {{.*}}Kernel3 + // 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: 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}} + AttrFunc(); + }); +} + +// The kernel has an attribute. +void calls_kernel_4() { + // 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}} + AttrFunc(); + }); +} 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/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/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..d139392620551 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size-device.cpp @@ -1,5 +1,5 @@ -// 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 -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" @@ -67,7 +67,7 @@ int main() { foo(); baz(); }); -#endif +#endif // TRIGGER_ERROR h.single_task([]() [[intel::reqd_sub_group_size(2)]]{}); h.single_task([]() [[intel::reqd_sub_group_size(4)]] { foo(); }); 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..19f22d391a72a 100644 --- a/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp +++ b/clang/test/SemaSYCL/scheduler_target_fmax_mhz_ast.cpp @@ -61,6 +61,7 @@ class KernelFunctor2 { int main() { deviceQueue.submit([&](sycl::handler &h) { + // Test attribute is propagated. // CHECK-LABEL: FunctionDecl {{.*}}kernel_name_1 // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr KernelFunctor f1; diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp new file mode 100644 index 0000000000000..5343019f43a08 --- /dev/null +++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp @@ -0,0 +1,41 @@ +// 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. + +#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]]{}); + + // 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(); }); + }); + 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.