From cf87963d73b3eb1d8049b25e72899be1a3cd5b65 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 9 Jul 2021 05:20:52 -0700 Subject: [PATCH 01/16] [SYCL] Implement SYCL 2020 spec functionality: no propagation from functions to the caller Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaSYCL.cpp | 38 ++- .../check-direct-attribute-propagation.cpp | 142 +++++++++ .../intel-fpga-no-global-work-offset.cpp | 2 +- .../CodeGenSYCL/intel-max-global-work-dim.cpp | 2 +- .../CodeGenSYCL/intel-max-work-group-size.cpp | 2 +- clang/test/CodeGenSYCL/intel-restrict.cpp | 2 +- .../test/CodeGenSYCL/num-simd-work-items.cpp | 2 +- .../test/CodeGenSYCL/reqd-sub-group-size.cpp | 2 +- .../test/CodeGenSYCL/reqd-work-group-size.cpp | 2 +- .../CodeGenSYCL/scheduler-target-fmax-mhz.cpp | 2 +- .../CodeGenSYCL/sycl-multi-kernel-attr.cpp | 2 +- .../check-direct-attribute-propagation.cpp | 279 ++++++++++++++++++ .../check-notdirect-attribute-propagation.cpp | 6 +- .../intel-fpga-no-global-work-offset.cpp | 2 +- .../intel-max-global-work-dim-device.cpp | 4 +- .../intel-max-work-group-size-device.cpp | 4 +- clang/test/SemaSYCL/intel-restrict.cpp | 4 +- clang/test/SemaSYCL/named_sub_group_size.cpp | 61 ++++ .../SemaSYCL/num_simd_work_items_device.cpp | 4 +- .../SemaSYCL/parallel_for_wrapper_attr.cpp | 2 +- .../redeclaration-attribute-propagation.cpp | 6 +- .../SemaSYCL/reqd-sub-group-size-device.cpp | 4 +- clang/test/SemaSYCL/sycl-esimd.cpp | 2 +- 23 files changed, 539 insertions(+), 37 deletions(-) create mode 100644 clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp create mode 100644 clang/test/SemaSYCL/check-direct-attribute-propagation.cpp create mode 100644 clang/test/SemaSYCL/named_sub_group_size.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 5126475f8aa4c..73237b01e0856 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -558,15 +558,35 @@ 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); + }); + } 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); + }); + } + } // Attributes that should not be propagated from device functions to a kernel. if (DirectlyCalled) { diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp new file mode 100644 index 0000000000000..59d973346db15 --- /dev/null +++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp @@ -0,0 +1,142 @@ +// 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" + +using namespace cl::sycl; +queue q; + +class Foo { +public: + [[intel::scheduler_target_fmax_mhz(1)]] void operator()() const {} +}; + +template +class Functor { +public: + [[intel::scheduler_target_fmax_mhz(SIZE)]] void operator()() const {} +}; + +[[intel::scheduler_target_fmax_mhz(5)]] void foo() {} + +class Foo1 { +public: + [[intel::num_simd_work_items(1)]] void operator()() const {} +}; + +template +class Functor1 { +public: + [[intel::num_simd_work_items(SIZE)]] void operator()() const {} +}; + +[[intel::num_simd_work_items(5)]] void foo1() {} + +class Foo2 { +public: + [[intel::no_global_work_offset(1)]] void operator()() const {} +}; + +template +class Functor2 { +public: + [[intel::no_global_work_offset(SIZE)]] void operator()() const {} +}; + +[[intel::no_global_work_offset(0)]] void foo2() {} + +class Foo3 { +public: + [[intel::max_global_work_dim(1)]] void operator()() const {} +}; + +template +class Functor3 { +public: + [[intel::max_global_work_dim(SIZE)]] void operator()() const {} +}; + +[[intel::max_global_work_dim(1)]] void foo3() {} + + +int main() { + q.submit([&](handler &h) { + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]] + Foo boo; + h.single_task(boo); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]] + h.single_task( + []() [[intel::scheduler_target_fmax_mhz(42)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]] + Functor<2> f; + h.single_task(f); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() + // CHECK-NOT: !scheduler_target_fmax_mhz + h.single_task( + []() { foo(); }); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !num_simd_work_items ![[NUM1]] + Foo1 boo1; + h.single_task(boo1); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 {{.*}} !num_simd_work_items ![[NUM42]] + h.single_task( + []() [[intel::num_simd_work_items(42)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 {{.*}} !num_simd_work_items ![[NUM2]] + Functor1<2> f1; + h.single_task(f1); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() + // CHECK-NOT: !num_simd_work_items + h.single_task( + []() { foo1(); }); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 {{.*}} !no_global_work_offset ![[NUM:[0-9]+]] + Foo2 boo2; + h.single_task(boo2); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name10() #0 {{.*}} ![[NUM0:[0-9]+]] + h.single_task( + []() [[intel::no_global_work_offset(0)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 {{.*}} !no_global_work_offset ![[NUM]] + Functor2<1> f2; + h.single_task(f2); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() + // CHECK-NOT: !no_global_work_offset + h.single_task( + []() { foo2(); }); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 {{.*}} !max_global_work_dim ![[NUM1]] + Foo3 boo3; + h.single_task(boo3); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 {{.*}} !max_global_work_dim ![[NUM1]] + h.single_task( + []() [[intel::max_global_work_dim(1)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 {{.*}} !max_global_work_dim ![[NUM2]] + Functor3<2> f3; + h.single_task(f3); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() + // CHECK-NOT: !max_global_work_dim + h.single_task( + []() { foo3(); }); + }); + return 0; +} + +// CHECK: ![[NUM]] = !{} +// CHECK: ![[NUM1]] = !{i32 1} +// CHECK: ![[NUM42]] = !{i32 42} +// CHECK: ![[NUM2]] = !{i32 2} +// CHECK-NOT: ![[NUM0]] = !{i32 0} 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 d85295bc53178..d63d6b076fd81 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,4 @@ -// 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 -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp index 53ab36ec9ab92..e3be51dd70b96 100644 --- a/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp +++ b/clang/test/CodeGenSYCL/intel-max-global-work-dim.cpp @@ -1,4 +1,4 @@ -// 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 -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp index 28c0bae228871..254a9175db744 100644 --- a/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/intel-max-work-group-size.cpp @@ -1,4 +1,4 @@ -// 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 -sycl-std=2017 -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/num-simd-work-items.cpp b/clang/test/CodeGenSYCL/num-simd-work-items.cpp index eebce1f408de4..56c7bd511e346 100644 --- a/clang/test/CodeGenSYCL/num-simd-work-items.cpp +++ b/clang/test/CodeGenSYCL/num-simd-work-items.cpp @@ -1,4 +1,4 @@ -// 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 -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp index 128a69052118e..01957696603e3 100644 --- a/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-sub-group-size.cpp @@ -1,4 +1,4 @@ -// 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 -disable-llvm-passes -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp index 0986e89738514..d9e79bd4c31af 100644 --- a/clang/test/CodeGenSYCL/reqd-work-group-size.cpp +++ b/clang/test/CodeGenSYCL/reqd-work-group-size.cpp @@ -1,4 +1,4 @@ -// 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 -sycl-std=2017 -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp index 55cf1f176e9fe..6bd6c0fbc7cf6 100644 --- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp +++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp @@ -1,4 +1,4 @@ -// 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 -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp index fd4e014d5825a..c316e611a18a8 100644 --- a/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp +++ b/clang/test/CodeGenSYCL/sycl-multi-kernel-attr.cpp @@ -1,4 +1,4 @@ -// 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 -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -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..131b587307b7f --- /dev/null +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -0,0 +1,279 @@ +// 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 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 { + [[sycl::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 {} +}; + +struct FuncObj7 { + [[intel::max_global_work_dim(1)]] 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() {} + +[[intel::max_global_work_dim(0)]] void func7() {} + +[[intel::reqd_sub_group_size(4)]] void func8() {} + +class Functor { +public: + void operator()() const { + func8(); + } +}; + +class Functor1 { +public: + [[intel::reqd_sub_group_size(12)]] void operator()() const {} +}; + +int main() { + deviceQueue.submit([&](sycl::handler &h) { + // CHECK: FunctionDecl {{.*}}test_kernel1 + // CHECK: SYCLSimdAttr {{.*}} Implicit + // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit + // CHECK-NEXT: SYCLSimdAttr + h.single_task( + FuncObj()); + // 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. + // 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. + // CHECK: FunctionDecl {{.*}}test_kernel4 + // CHECK-NOT: SYCLIntelNoGlobalWorkOffsetAttr + h.single_task( + []() { func1(); }); + + // CHECK: FunctionDecl {{.*}}test_kernel5 + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task( + FuncObj1()); + + // 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]]{}); + + // CHECK: FunctionDecl {{.*}}test_kernel7 + // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 10 + // CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} + h.single_task( + FuncObj2()); + + // 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. + // CHECK: FunctionDecl {{.*}}test_kernel9 + // CHECK-NOT: SYCLIntelSchedulerTargetFmaxMhzAttr + h.single_task( + []() { func2(); }); + + // 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. + // CHECK: FunctionDecl {{.*}}test_kernel11 + // CHECK-NOT: SYCLIntelMaxWorkGroupSizeAttr + h.single_task( + []() { func3(); }); + + // 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)]]{}); + + // 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. + // CHECK: FunctionDecl {{.*}}test_kernel14 + // CHECK-NOT: ReqdWorkGroupSizeAttr + h.single_task( + []() { func4(); }); + + // 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)]]{}); + + // CHECK: FunctionDecl {{.*}}test_kernel16 + // CHECK: SYCLIntelNumSimdWorkItemsAttr + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + h.single_task( + FuncObj5()); + + // 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. + // CHECK: FunctionDecl {{.*}}test_kernel18 + // CHECK-NOT: SYCLIntelNumSimdWorkItemsAttr + h.single_task( + []() { func5(); }); + + // CHECK: FunctionDecl {{.*}}test_kernel19 + // CHECK: SYCLIntelKernelArgsRestrictAttr + h.single_task( + FuncObj6()); + + // CHECK: FunctionDecl {{.*}}test_kernel20 + // CHECK: SYCLIntelKernelArgsRestrictAttr + h.single_task( + []() [[intel::kernel_args_restrict]]{}); + + // Test attribute is not propagated. + // CHECK: FunctionDecl {{.*}}test_kernel21 + // CHECK-NOT: SYCLIntelKernelArgsRestrictAttr + h.single_task( + []() { func6(); }); + + // CHECK: FunctionDecl {{.*}}test_kernel22 + // CHECK: SYCLIntelMaxGlobalWorkDimAttr + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + h.single_task( + FuncObj7()); + + // CHECK: FunctionDecl {{.*}}test_kernel23 + // CHECK: SYCLIntelMaxGlobalWorkDimAttr + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 0 + // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + h.single_task( + []() [[intel::max_global_work_dim(0)]]{}); + + // Test attribute is not propagated. + // CHECK: FunctionDecl {{.*}}test_kernel24 + // CHECK-NOT: SYCLIntelMaxGlobalWorkDimAttr + h.single_task( + []() { func7(); }); + + // Test attribute is not propagated. + // CHECK: FunctionDecl {{.*}}test_kernel25 + // CHECK-NOT: IntelReqdSubGroupSizeAttr + Functor f; + h.single_task(f); + + // CHECK: FunctionDecl {{.*}}test_kernel26 + // CHECK: IntelReqdSubGroupSizeAttr + Functor1 f1; + h.single_task(f1); + }); + 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/intel-fpga-no-global-work-offset.cpp b/clang/test/SemaSYCL/intel-fpga-no-global-work-offset.cpp index b3c2d4b8e40c2..d85ca9ddc1f96 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" 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 a93096acec29e..507001a889432 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=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-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-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/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 0841662f9b7bb..858fe713e412b 100644 --- a/clang/test/SemaSYCL/num_simd_work_items_device.cpp +++ b/clang/test/SemaSYCL/num_simd_work_items_device.cpp @@ -1,5 +1,5 @@ -// 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 -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -sycl-std=2020 -Wno-sycl-2017-compat -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsycl-is-device -internal-isystem %S/Inputs -triple spir64 -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp b/clang/test/SemaSYCL/parallel_for_wrapper_attr.cpp index dfc39bc84a6d2..6be07189d4d1d 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 -ast-dump -sycl-std=2017 -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..4569ca5284e83 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" 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 f644e4578bfea840e349832b591e583be9b4bc03 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 9 Jul 2021 06:15:18 -0700 Subject: [PATCH 02/16] add test Signed-off-by: Soumi Manna --- .../check-direct-attribute-propagation.cpp | 118 ++++++++++++++++++ 1 file changed, 118 insertions(+) diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp index 59d973346db15..055735f0175a0 100644 --- a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp @@ -58,6 +58,66 @@ class Functor3 { [[intel::max_global_work_dim(1)]] void foo3() {} +class Foo4 { +public: + [[intel::reqd_sub_group_size(16)]] void operator()() const {} +}; + +[[intel::reqd_sub_group_size(8)]] void foo4() {} + +class Functor4 { +public: + void operator()() const { + foo(); + } +}; + +template +class Functor5 { +public: + [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {} +}; + +class Foo5 { +public: + [[sycl::reqd_work_group_size(32, 16, 16)]] void operator()() const {} +}; + +[[sycl::reqd_work_group_size(8, 1, 1)]] void foo5() {} + +class Functor6 { +public: + void operator()() const { + foo5(); + } +}; + +template +class Functor7 { +public: + [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} +}; + +class Foo6 { +public: + [[intel::max_work_group_size(32, 16, 16)]] void operator()() const {} +}; + +[[intel::max_work_group_size(8, 1, 1)]] void foo6() {} + +class Functor8 { +public: + void operator()() const { + foo6(); + } +}; + +template +class Functor9 { +public: + [[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} +}; + int main() { q.submit([&](handler &h) { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]] @@ -131,6 +191,60 @@ int main() { // CHECK-NOT: !max_global_work_dim h.single_task( []() { foo3(); }); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM16:[0-9]+]] + Foo4 boo4; + h.single_task(boo4); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM1]] + h.single_task( + []() [[intel::reqd_sub_group_size(1)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM2]] + Functor5<2> f5; + h.single_task(f5); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() + // CHECK-NOT: !reqd_sub_group_size + Functor4 f4; + h.single_task(f4); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 {{.*}} !reqd_work_group_size ![[NUM32:[0-9]+]] + Foo5 boo5; + h.single_task(boo5); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 {{.*}} !reqd_work_group_size ![[NUM88:[0-9]+]] + h.single_task( + []() [[sycl::reqd_work_group_size(8, 8, 8)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 {{.*}} !reqd_work_group_size ![[NUM22:[0-9]+]] + Functor7<2, 2, 2> f7; + h.single_task(f7); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() + // CHECK-NOT: !reqd_work_group_size + Functor6 f6; + h.single_task(f6); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 {{.*}} !max_work_group_size ![[NUM32]] + Foo6 boo6; + h.single_task(boo6); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 {{.*}} !max_work_group_size ![[NUM88]] + h.single_task( + []() [[intel::max_work_group_size(8, 8, 8)]]{}); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 {{.*}} !max_work_group_size ![[NUM22]] + Functor9<2, 2, 2> f9; + h.single_task(f9); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() + // CHECK-NOT: !max_work_group_size + Functor8 f8; + h.single_task(f8); }); return 0; } @@ -140,3 +254,7 @@ int main() { // CHECK: ![[NUM42]] = !{i32 42} // CHECK: ![[NUM2]] = !{i32 2} // CHECK-NOT: ![[NUM0]] = !{i32 0} +// CHECK: ![[NUM16]] = !{i32 16} +// CHECK: ![[NUM32]] = !{i32 16, i32 16, i32 32} +// CHECK: ![[NUM88]] = !{i32 8, i32 8, i32 8} +// CHECK: ![[NUM22]] = !{i32 2, i32 2, i32 2} From 351d0c439f1c125911be810e49712ab3e9a48b84 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Fri, 9 Jul 2021 08:03:03 -0700 Subject: [PATCH 03/16] Update tests Signed-off-by: Soumi Manna --- .../check-direct-attribute-propagation.cpp | 13 +- .../check-direct-attribute-propagation.cpp | 20 ++- ...eqd-work-group-size-device-direct-prop.cpp | 146 ++++++++++++++++++ .../intel-reqd-work-group-size-device.cpp | 21 +-- ...eqd-work-group-size-device-direct-prop.cpp | 116 ++++++++++++++ .../SemaSYCL/reqd-work-group-size-device.cpp | 42 +---- 6 files changed, 296 insertions(+), 62 deletions(-) create mode 100644 clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp create mode 100644 clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp index 055735f0175a0..7c91aea0c1154 100644 --- a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp @@ -1,5 +1,10 @@ // 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 +// Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], +// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], +// [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], and +// [[intel::max_work_group_size()]] kernel attributes in SYCL 2020. + #include "sycl.hpp" using namespace cl::sycl; @@ -60,10 +65,10 @@ class Functor3 { class Foo4 { public: - [[intel::reqd_sub_group_size(16)]] void operator()() const {} + [[sycl::reqd_sub_group_size(16)]] void operator()() const {} }; -[[intel::reqd_sub_group_size(8)]] void foo4() {} +[[sycl::reqd_sub_group_size(8)]] void foo4() {} class Functor4 { public: @@ -75,7 +80,7 @@ class Functor4 { template class Functor5 { public: - [[intel::reqd_sub_group_size(SIZE)]] void operator()() const {} + [[sycl::reqd_sub_group_size(SIZE)]] void operator()() const {} }; class Foo5 { @@ -198,7 +203,7 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM1]] h.single_task( - []() [[intel::reqd_sub_group_size(1)]]{}); + []() [[sycl::reqd_sub_group_size(1)]]{}); // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM2]] Functor5<2> f5; diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp index 131b587307b7f..2c762838a1dbd 100644 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -1,6 +1,9 @@ // 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 kernel attributes to the caller when used on a function. +// Tests for AST of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], +// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], [[intel::sycl_explicit_simd]], +// [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], and +// [[intel::max_work_group_size()]] function attributes in SYCL 2020. #include "sycl.hpp" @@ -46,7 +49,7 @@ struct FuncObj7 { [[intel::max_work_group_size(1, 1, 1)]] void func3() {} -[[intel::reqd_work_group_size(1, 1, 1)]] void func4() {} +[[sycl::reqd_work_group_size(1, 1, 1)]] void func4() {} [[intel::num_simd_work_items(5)]] void func5() {} @@ -54,7 +57,7 @@ struct FuncObj7 { [[intel::max_global_work_dim(0)]] void func7() {} -[[intel::reqd_sub_group_size(4)]] void func8() {} +[[sycl::reqd_sub_group_size(4)]] void func8() {} class Functor { public: @@ -65,7 +68,7 @@ class Functor { class Functor1 { public: - [[intel::reqd_sub_group_size(12)]] void operator()() const {} + [[sycl::reqd_sub_group_size(12)]] void operator()() const {} }; int main() { @@ -202,7 +205,7 @@ int main() { // CHECK-NEXT: value: Int 8 // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} h.single_task( - []() [[intel::reqd_work_group_size(8, 8, 8)]]{}); + []() [[sycl::reqd_work_group_size(8, 8, 8)]]{}); // CHECK: FunctionDecl {{.*}}test_kernel16 // CHECK: SYCLIntelNumSimdWorkItemsAttr @@ -274,6 +277,13 @@ int main() { // CHECK: IntelReqdSubGroupSizeAttr Functor1 f1; h.single_task(f1); + + // CHECK: FunctionDecl {{.*}}test_kernel27 + // CHECK: IntelReqdSubGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 8 + h.single_task( + []() [[sycl::reqd_sub_group_size(8)]]{}); }); return 0; } diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp new file mode 100644 index 0000000000000..ba814f82ef025 --- /dev/null +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp @@ -0,0 +1,146 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -fsyntax-only -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s + +// Test for AST of reqd_work_group_size kernel attribute in SYCL 2020. + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; + +#ifndef __SYCL_DEVICE_ONLY__ +// expected-no-diagnostics +class Functor { +public: + [[sycl::reqd_work_group_size(4)]] void operator()() const {} +}; + +void bar() { + q.submit([&](handler &h) { + Functor f; + h.single_task(f); + }); +} + +#else +#ifdef TRIGGER_ERROR +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 // TRIGGER_ERROR + +class Functor33 { +public: + // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + [[sycl::reqd_work_group_size(32, -4)]] void operator()() const {} +}; + +class Functor30 { +public: + // expected-warning@+1 2{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} + [[sycl::reqd_work_group_size(30, -30, -30)]] void operator()() const {} +}; + +class Functor16 { +public: + [[sycl::reqd_work_group_size(16)]] void operator()() const {} +}; + +class Functor64 { +public: + [[sycl::reqd_work_group_size(64, 64)]] void operator()() const {} +}; + +class Functor16x16x16 { +public: + [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \ + // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} +}; + +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?}} +}; + +int main() { + q.submit([&](handler &h) { + Functor16 f16; + h.single_task(f16); + + Functor16x16x16 f16x16x16; + h.single_task(f16x16x16); + + FunctorAttr fattr; + h.single_task(fattr); + + Functor33 f33; + h.single_task(f33); + + Functor30 f30; + h.single_task(f30); + }); + return 0; +} + +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 16 +// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 +// 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{{$}} +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 128 +// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 128 +// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 128 +// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 32 +// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int -4 +// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' +// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 30 +// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int -30 +// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' +// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int -30 +// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' +#endif // __SYCL_DEVICE_ONLY__ 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 52c9cdfe9a6f5..e06ae8bd70d6b 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 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -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 + +// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1. #include "sycl.hpp" @@ -30,15 +32,6 @@ void bar() { [[sycl::reqd_work_group_size(32, 32)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} [[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} -#ifdef TRIGGER_ERROR -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 // TRIGGER_ERROR - class Functor33 { public: // expected-warning@+1{{implicit conversion changes signedness: 'int' to 'unsigned long long'}} @@ -63,8 +56,7 @@ class Functor64 { class Functor16x16x16 { public: - [[intel::reqd_work_group_size(16, 16, 16)]] void operator()() const {} // expected-warning {{attribute 'intel::reqd_work_group_size' is deprecated}} \ - // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} + [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} }; class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} @@ -83,8 +75,7 @@ class Functor { 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?}} + [[sycl::reqd_work_group_size(128, 128, 128)]] void operator()() const {} }; int main() { diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp new file mode 100644 index 0000000000000..173406394d49c --- /dev/null +++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp @@ -0,0 +1,116 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump %s | FileCheck %s + +// Test for AST of reqd_work_group_size kernel attribute in SYCL 2020. + +#include "sycl.hpp" + +using namespace cl::sycl; +queue q; + +// Same for the default values. +// FIXME: This turns out to be wrong as there aren't really default values +// (that is an implementation detail we use but shouldn't expose to the user). +// Instead, the dimensionality of the attribute needs to match that of the +// kernel, so the one, two, and three arg forms of the attribute are actually +// *different* attributes. This means that you should not be able to redeclare +// the function with a different dimensionality. +[[sycl::reqd_work_group_size(4)]] void four_again(); +[[sycl::reqd_work_group_size(4)]] void four_again(); // OK +[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK +[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK +[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK +[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK + +// The GNU and [[cl::reqd_work_group_size]] spellings are deprecated in SYCL +// mode, and still requires all three arguments. +__attribute__((reqd_work_group_size(4, 4, 4))) void four_once_more(); // expected-warning {{attribute '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, 4, 4)]] void four_with_feeling(); // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ + // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} + +#ifdef TRIGGER_ERROR +__attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \ + // expected-warning {{attribute '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)]] void four_with_more_feeling(); // 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?}} + +// Make sure there's at least one argument passed for the SYCL spelling. +[[sycl::reqd_work_group_size]] void four_no_more(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}} +#endif // TRIGGER_ERROR + +class Functor16 { +public: + [[sycl::reqd_work_group_size(16, 1, 1)]] [[sycl::reqd_work_group_size(16, 1, 1)]] void operator()() const {} +}; + +#ifdef TRIGGER_ERROR +class Functor32 { +public: + // expected-note@+3{{conflicting attribute is here}} + // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}} + // 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 +class Functor16x16x16 { +public: + [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} +}; + +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?}} +}; + +int main() { + q.submit([&](handler &h) { + Functor16 f16; + h.single_task(f16); + + Functor16x16x16 f16x16x16; + h.single_task(f16x16x16); + + FunctorAttr fattr; + h.single_task(fattr); + }); + return 0; +} + +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 16 +// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 +// 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{{$}} +// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 +// CHECK: ReqdWorkGroupSizeAttr {{.*}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 128 +// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 128 +// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 128 +// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp index cda6a368be58c..700f861c7fce9 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp @@ -1,6 +1,7 @@ -// 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 +// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1. #include "sycl.hpp" using namespace cl::sycl; @@ -34,23 +35,8 @@ queue q; [[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK [[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK -// The GNU and [[cl::reqd_work_group_size]] spellings are deprecated in SYCL -// mode, and still requires all three arguments. -__attribute__((reqd_work_group_size(4, 4, 4))) void four_once_more(); // expected-warning {{attribute '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, 4, 4)]] void four_with_feeling(); // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ - // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} - -#ifdef TRIGGER_ERROR -__attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-error {{'reqd_work_group_size' attribute requires exactly 3 arguments}} \ - // expected-warning {{attribute '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)]] void four_with_more_feeling(); // 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?}} - // Make sure there's at least one argument passed for the SYCL spelling. +#ifdef TRIGGER_ERROR [[sycl::reqd_work_group_size]] void four_no_more(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}} #endif // TRIGGER_ERROR @@ -87,12 +73,6 @@ class Functor { } }; -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?}} -}; - int main() { q.submit([&](handler &h) { Functor16 f16; @@ -104,9 +84,6 @@ int main() { Functor16x16x16 f16x16x16; h.single_task(f16x16x16); - FunctorAttr fattr; - h.single_task(fattr); - h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32), sycl::reqd_work_group_size(32, 32, 32)]] { f32x32x32(); }); @@ -176,17 +153,6 @@ int main() { // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 16 // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 // CHECK: ReqdWorkGroupSizeAttr {{.*}} // CHECK-NEXT: ConstantExpr{{.*}}'int' From a632cdec48c4811a01e78171100a632970c6ceab Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 12 Jul 2021 10:06:29 -0700 Subject: [PATCH 04/16] update doc Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 44 ++++++++++++++++++++++----- 1 file changed, 36 insertions(+), 8 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 9320e56e69c6f..eb9bafe80ff93 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++ @@ -2656,6 +2663,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 +``cl::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() {} @@ -2800,8 +2812,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++ @@ -2832,8 +2846,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++ @@ -2890,6 +2906,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() {} @@ -2920,6 +2940,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]] @@ -4607,6 +4631,10 @@ the ``[[intel::named_sub_group_size(NAME)]]`` documentation for clarification. This attribute is mutually exclusive with ``[[intel::named_sub_group_size(NAME)]]`` and ``[[intel::sycl_explicit_simd]]``. +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. + 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: From 48d83969a9071902e653445cc98b63de2716ff83 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 12 Jul 2021 15:50:18 -0700 Subject: [PATCH 05/16] remove duplicate ones Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaSYCL.cpp | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 73237b01e0856..82f20295d9c9c 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -576,14 +576,12 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, // 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); + return isa(A); }); } } From 828dce2e830bba72dfbe918e9ab211641201f7cd Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 12 Jul 2021 15:55:39 -0700 Subject: [PATCH 06/16] fix clang format issues Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaSYCL.cpp | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 82f20295d9c9c..8d937bb77ac73 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -577,11 +577,11 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, if (DirectlyCalled) { llvm::copy_if(FD->getAttrs(), std::back_inserter(Attrs), [](Attr *A) { return isa(A); + ReqdWorkGroupSizeAttr, SYCLIntelKernelArgsRestrictAttr, + SYCLIntelNumSimdWorkItemsAttr, + SYCLIntelSchedulerTargetFmaxMhzAttr, + SYCLIntelMaxWorkGroupSizeAttr, SYCLIntelMaxGlobalWorkDimAttr, + SYCLIntelNoGlobalWorkOffsetAttr, SYCLSimdAttr>(A); }); } } From 1c76c147a95c6b89fc9b0461c062d7dffca25a45 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Mon, 12 Jul 2021 20:31:46 -0700 Subject: [PATCH 07/16] update tests Signed-off-by: Soumi Manna --- .../check-direct-attribute-propagation.cpp | 60 ++++++++++++++++++- clang/test/SemaSYCL/named_sub_group_size.cpp | 2 - 2 files changed, 57 insertions(+), 5 deletions(-) diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp index 7c91aea0c1154..c6cffa9aab964 100644 --- a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp @@ -1,9 +1,9 @@ // 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 // Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], -// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], -// [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], and -// [[intel::max_work_group_size()]] kernel attributes in SYCL 2020. +// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], [[sycl::reqd_sub_group_size()]], +// [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], [[intel::max_work_group_size()]], +// and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020. #include "sycl.hpp" @@ -123,6 +123,27 @@ class Functor9 { [[intel::max_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() const {} }; +class Foo7 { +public: + [[intel::sycl_explicit_simd]] void operator()() const {} +}; + +[[intel::sycl_explicit_simd]] void foo7() {} + +class Foo8 { +public: + [[intel::kernel_args_restrict]] void operator()() const {} +}; + +[[intel::kernel_args_restrict]] void foo8() {} + +class Functor10 { +public: + void operator()() const { + foo8(); + } +}; + int main() { q.submit([&](handler &h) { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]] @@ -250,6 +271,39 @@ int main() { // CHECK-NOT: !max_work_group_size Functor8 f8; h.single_task(f8); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() + // CHECK-NOT: !sycl_explicit_simd + // CHECK-SAME: { + // CHECK: define {{.*}}spir_func void @{{.*}}foo7{{.*}} !sycl_explicit_simd ![[NUM]] + h.single_task( + []() { foo7(); }); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name30() #0 !intel_reqd_sub_group_size ![[NUM1]] !sycl_explicit_simd ![[NUM]] + Foo7 boo7; + h.single_task(boo7); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name31() #0 !intel_reqd_sub_group_size ![[NUM1]] !sycl_explicit_simd ![[NUM]] + h.single_task( + []() [[intel::sycl_explicit_simd]]{}); + + // Test attribute is not propagated. + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0 !kernel_arg_buffer_location ![[NUM]] + // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class._ZTS9Functor10.Functor10 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 + // CHECK-NOT: noalias + Functor10 f10; + h.single_task(f10); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name33() #0 !kernel_arg_buffer_location ![[NUM]] + // CHECK: define {{.*}}spir_func void @{{.*}}Foo8{{.*}}(%class.{{.*}}Foo8.Foo8 addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 + Foo8 boo8; + h.single_task(boo8); + + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name34() #0 !kernel_arg_buffer_location ![[NUM]] + // CHECK: define {{.*}}spir_func void @{{.*}}(%class.{{.*}}.anon addrspace(4)* noalias align 1 dereferenceable_or_null(1) %this) #4 align 2 + h.single_task( + []() [[intel::kernel_args_restrict]]{}); }); return 0; } diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index 13ce54c4f0116..981a9fb82d851 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -5,7 +5,6 @@ #include "Inputs/sycl.hpp" -// The kernel has an attribute. void calls_kernel_1() { // CHECK: FunctionDecl {{.*}}Kernel1 // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic @@ -49,7 +48,6 @@ void calls_kernel_3() { }); } -// The kernel has an attribute. void calls_kernel_4() { // CHECK: FunctionDecl {{.*}}Kernel5 // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic From fcc5de0182ca3f86dc0515fcb18b41831f90061e Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 13 Jul 2021 13:54:35 -0700 Subject: [PATCH 08/16] address @aaron's review comments Signed-off-by: Soumi Manna --- clang/include/clang/Basic/AttrDocs.td | 8 ++++---- clang/lib/Sema/SemaSYCL.cpp | 22 +++++----------------- 2 files changed, 9 insertions(+), 21 deletions(-) diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index eb9bafe80ff93..f2ed088e268a0 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -2663,10 +2663,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. -In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size`` or -``cl::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. +In SYCL 1.2.1 mode, the ``intel::reqd_work_group_size``, +``cl::reqd_work_group_size``, and ``sycl::reqd_work_group_size`` attributes are +propagated from the function they are applied to onto the kernel which calls the +function. In SYCL 2020 mode, the attributes are not propagated to the kernel. .. code-block:: c++ diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 8d937bb77ac73..6104069f75f11 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -558,10 +558,11 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, if (!FD->hasAttrs()) return; - // Attributes that should be propagated from device functions to a kernel - // in SYCL 1.2.1. - if (S.getASTContext().getLangOpts().getSYCLVersion() < - LangOptions::SYCL_2020) { + // In SYCL 1.2.1 mode, the attributes are propagated from the function they + // are applied to onto the kernel which calls the function. + // In SYCL 2020 mode, the attributes are not propagated to the kernel. + if (DirectlyCalled || + 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); }); - } 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(A); - }); - } } // Attributes that should not be propagated from device functions to a kernel. From def5f8f02e0d5d3c7d347906333a14c9f553e00e Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Tue, 13 Jul 2021 14:01:13 -0700 Subject: [PATCH 09/16] fix format issues Signed-off-by: Soumi Manna --- clang/lib/Sema/SemaSYCL.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index 6104069f75f11..acaefa9f91a36 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -561,8 +561,8 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, // In SYCL 1.2.1 mode, the attributes are propagated from the function they // are applied to onto the kernel which calls the function. // In SYCL 2020 mode, the attributes are not propagated to the kernel. - if (DirectlyCalled || - S.getASTContext().getLangOpts().getSYCLVersion() < LangOptions::SYCL_2020) { + if (DirectlyCalled || 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 Date: Wed, 14 Jul 2021 13:41:29 -0700 Subject: [PATCH 10/16] address @elizabeth's review comments Signed-off-by: Soumi Manna --- .../check-direct-attribute-propagation.cpp | 76 +++++++++++-------- .../check-direct-attribute-propagation.cpp | 6 +- clang/test/SemaSYCL/sycl-esimd-ast.cpp | 41 ++++++++++ 3 files changed, 90 insertions(+), 33 deletions(-) create mode 100644 clang/test/SemaSYCL/sycl-esimd-ast.cpp diff --git a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp index c6cffa9aab964..a04610d556360 100644 --- a/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/CodeGenSYCL/check-direct-attribute-propagation.cpp @@ -1,7 +1,7 @@ // 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 // Tests for IR of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], -// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], [[sycl::reqd_sub_group_size()]], +// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[sycl::reqd_sub_group_size()]], // [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], [[intel::max_work_group_size()]], // and [[intel::sycl_explicit_simd]] function attributes in SYCL 2020. @@ -73,7 +73,7 @@ class Foo4 { class Functor4 { public: void operator()() const { - foo(); + foo4(); } }; @@ -146,43 +146,47 @@ class Functor10 { int main() { q.submit([&](handler &h) { - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name1() #0 !kernel_arg_buffer_location ![[NUM:[0-9]+]] !scheduler_target_fmax_mhz ![[NUM1:[0-9]+]] Foo boo; h.single_task(boo); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name2() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM42:[0-9]+]] h.single_task( []() [[intel::scheduler_target_fmax_mhz(42)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 {{.*}} !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name3() #0 !kernel_arg_buffer_location ![[NUM]] !scheduler_target_fmax_mhz ![[NUM2:[0-9]+]] Functor<2> f; h.single_task(f); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name4() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !scheduler_target_fmax_mhz + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z3foov() h.single_task( []() { foo(); }); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 {{.*}} !num_simd_work_items ![[NUM1]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name5() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM1]] Foo1 boo1; h.single_task(boo1); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 {{.*}} !num_simd_work_items ![[NUM42]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name6() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM42]] h.single_task( []() [[intel::num_simd_work_items(42)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 {{.*}} !num_simd_work_items ![[NUM2]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name7() #0 !kernel_arg_buffer_location ![[NUM]] !num_simd_work_items ![[NUM2]] Functor1<2> f1; h.single_task(f1); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name8() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !num_simd_work_items + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo1v() h.single_task( []() { foo1(); }); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 {{.*}} !no_global_work_offset ![[NUM:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name9() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM:[0-9]+]] Foo2 boo2; h.single_task(boo2); @@ -190,90 +194,100 @@ int main() { h.single_task( []() [[intel::no_global_work_offset(0)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 {{.*}} !no_global_work_offset ![[NUM]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name11() #0 !kernel_arg_buffer_location ![[NUM]] !no_global_work_offset ![[NUM]] Functor2<1> f2; h.single_task(f2); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name12() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !no_global_work_offset + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo2v() h.single_task( []() { foo2(); }); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 {{.*}} !max_global_work_dim ![[NUM1]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name13() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]] Foo3 boo3; h.single_task(boo3); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 {{.*}} !max_global_work_dim ![[NUM1]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name14() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM1]] h.single_task( []() [[intel::max_global_work_dim(1)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 {{.*}} !max_global_work_dim ![[NUM2]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name15() #0 !kernel_arg_buffer_location ![[NUM]] !max_global_work_dim ![[NUM2]] Functor3<2> f3; h.single_task(f3); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name16() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !max_global_work_dim + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo3v() h.single_task( []() { foo3(); }); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM16:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name17() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM16:[0-9]+]] Foo4 boo4; h.single_task(boo4); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM1]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name18() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM1]] h.single_task( []() [[sycl::reqd_sub_group_size(1)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 {{.*}} !intel_reqd_sub_group_size ![[NUM2]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name19() #0 !kernel_arg_buffer_location ![[NUM]] !intel_reqd_sub_group_size ![[NUM2]] Functor5<2> f5; h.single_task(f5); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name20() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !reqd_sub_group_size + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo4v() Functor4 f4; h.single_task(f4); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 {{.*}} !reqd_work_group_size ![[NUM32:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name21() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM32:[0-9]+]] Foo5 boo5; h.single_task(boo5); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 {{.*}} !reqd_work_group_size ![[NUM88:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name22() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM88:[0-9]+]] h.single_task( []() [[sycl::reqd_work_group_size(8, 8, 8)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 {{.*}} !reqd_work_group_size ![[NUM22:[0-9]+]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name23() #0 !kernel_arg_buffer_location ![[NUM]] !reqd_work_group_size ![[NUM22:[0-9]+]] Functor7<2, 2, 2> f7; h.single_task(f7); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name24() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !reqd_work_group_size + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo5v() Functor6 f6; h.single_task(f6); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 {{.*}} !max_work_group_size ![[NUM32]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name25() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM32]] Foo6 boo6; h.single_task(boo6); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 {{.*}} !max_work_group_size ![[NUM88]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name26() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM88]] h.single_task( []() [[intel::max_work_group_size(8, 8, 8)]]{}); - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 {{.*}} !max_work_group_size ![[NUM22]] + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name27() #0 !kernel_arg_buffer_location ![[NUM]] !max_work_group_size ![[NUM22]] Functor9<2, 2, 2> f9; h.single_task(f9); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name28() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !max_work_group_size + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo6v() Functor8 f8; h.single_task(f8); // Test attribute is not propagated. - // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() + // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name29() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK-NOT: !sycl_explicit_simd // CHECK-SAME: { // CHECK: define {{.*}}spir_func void @{{.*}}foo7{{.*}} !sycl_explicit_simd ![[NUM]] @@ -292,6 +306,8 @@ int main() { // CHECK: define {{.*}}spir_kernel void @{{.*}}kernel_name32() #0 !kernel_arg_buffer_location ![[NUM]] // CHECK: define {{.*}}spir_func void @{{.*}}Functor10{{.*}}(%class._ZTS9Functor10.Functor10 addrspace(4)* align 1 dereferenceable_or_null(1) %this) #3 comdat align 2 // CHECK-NOT: noalias + // CHECK-SAME: { + // CHECK: define dso_local spir_func void @_Z4foo8v() Functor10 f10; h.single_task(f10); diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp index 2c762838a1dbd..95941780fd9bc 100644 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -1,7 +1,7 @@ // RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -ast-dump %s | FileCheck %s // Tests for AST of [[intel::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], -// [[intel::no_global_work_offset()]], [[intel::no_global_work_offset()]], [[intel::sycl_explicit_simd]], +// [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[intel::sycl_explicit_simd]], // [[sycl::reqd_sub_group_size()]], [[sycl::reqd_work_group_size()]], [[intel::kernel_args_restrict]], and // [[intel::max_work_group_size()]] function attributes in SYCL 2020. @@ -109,12 +109,12 @@ int main() { h.single_task( FuncObj1()); - // CHECK: FunctionDecl {{.*}}test_kerne6 + // CHECK: FunctionDecl {{.*}}test_kernel6 // CHECK: SYCLIntelNoGlobalWorkOffsetAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - h.single_task( + h.single_task( []() [[intel::no_global_work_offset]]{}); // CHECK: FunctionDecl {{.*}}test_kernel7 diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp new file mode 100644 index 0000000000000..86454f7f06724 --- /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 in SYCL 1.2.1. + +#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; +} From 2b90071c4a5d153dca60f2ed94a67921bf09a288 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 14 Jul 2021 17:40:21 -0700 Subject: [PATCH 11/16] update tests Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/named_sub_group_size.cpp | 26 ++--- ...eqd-work-group-size-device-direct-prop.cpp | 95 +++++-------------- 2 files changed, 32 insertions(+), 89 deletions(-) diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index 981a9fb82d851..35552d282fc9c 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -5,13 +5,6 @@ #include "Inputs/sycl.hpp" -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 { } @@ -23,6 +16,13 @@ struct Functor1 { }; // Test attributes get propgated to the kernel. +void calls_kernel_1() { + // CHECK: FunctionDecl {{.*}}Kernel1 + // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic + sycl::kernel_single_task([]() [[intel::named_sub_group_size(automatic)]] { + }); +} + void calls_kernel_2() { Functor F; // CHECK: FunctionDecl {{.*}}Kernel2 @@ -35,9 +35,9 @@ void calls_kernel_2() { sycl::kernel_single_task(F1); } -// Test ttribute does not get propgated to the kernel. [[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc +// Test ttribute does not get propgated to the kernel. void calls_kernel_3() { // CHECK: FunctionDecl {{.*}}Kernel4 // CHECK-NOT: IntelNamedSubGroupSizeAttr {{.*}} @@ -47,13 +47,3 @@ void calls_kernel_3() { AttrFunc(); }); } - -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/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp index 173406394d49c..daff77d960106 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp @@ -8,20 +8,6 @@ using namespace cl::sycl; queue q; -// Same for the default values. -// FIXME: This turns out to be wrong as there aren't really default values -// (that is an implementation detail we use but shouldn't expose to the user). -// Instead, the dimensionality of the attribute needs to match that of the -// kernel, so the one, two, and three arg forms of the attribute are actually -// *different* attributes. This means that you should not be able to redeclare -// the function with a different dimensionality. -[[sycl::reqd_work_group_size(4)]] void four_again(); -[[sycl::reqd_work_group_size(4)]] void four_again(); // OK -[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK -[[sycl::reqd_work_group_size(4, 1)]] void four_again(); // OK -[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK -[[sycl::reqd_work_group_size(4, 1, 1)]] void four_again(); // OK - // The GNU and [[cl::reqd_work_group_size]] spellings are deprecated in SYCL // mode, and still requires all three arguments. __attribute__((reqd_work_group_size(4, 4, 4))) void four_once_more(); // expected-warning {{attribute 'reqd_work_group_size' is deprecated}} \ @@ -38,24 +24,6 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \ // expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} -// Make sure there's at least one argument passed for the SYCL spelling. -[[sycl::reqd_work_group_size]] void four_no_more(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}} -#endif // TRIGGER_ERROR - -class Functor16 { -public: - [[sycl::reqd_work_group_size(16, 1, 1)]] [[sycl::reqd_work_group_size(16, 1, 1)]] void operator()() const {} -}; - -#ifdef TRIGGER_ERROR -class Functor32 { -public: - // expected-note@+3{{conflicting attribute is here}} - // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}} - // 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 class Functor16x16x16 { public: [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} @@ -69,48 +37,33 @@ class FunctorAttr { int main() { q.submit([&](handler &h) { - Functor16 f16; - h.single_task(f16); - + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 + // 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{{$}} Functor16x16x16 f16x16x16; - h.single_task(f16x16x16); + h.single_task(f16x16x16); + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 + // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 128 + // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 128 + // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 128 + // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} FunctorAttr fattr; - h.single_task(fattr); + h.single_task(fattr); }); return 0; } - -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 -// 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{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} From 960d52cac67cb98151cf665958b5924b4d3c433e Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Wed, 14 Jul 2021 18:47:34 -0700 Subject: [PATCH 12/16] fix test failures Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/named_sub_group_size.cpp | 10 ++++++++++ .../reqd-work-group-size-device-direct-prop.cpp | 12 ++++++++++++ 2 files changed, 22 insertions(+) diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index 35552d282fc9c..81972804a69a4 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -47,3 +47,13 @@ void calls_kernel_3() { AttrFunc(); }); } + +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/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp index daff77d960106..a350f344f00f2 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp @@ -23,6 +23,18 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro [[cl::reqd_work_group_size(4)]] void four_with_more_feeling(); // 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 // TRIGGER_ERROR + + +#ifdef TRIGGER_ERROR +class Functor32 { +public: + // expected-note@+3{{conflicting attribute is here}} + // expected-warning@+2{{attribute 'reqd_work_group_size' is already applied with different arguments}} + // 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 // TRIGGER_ERROR class Functor16x16x16 { public: From 40b69668ef6e45c817c3157df4add74b1b88edcf Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 15 Jul 2021 04:50:17 -0700 Subject: [PATCH 13/16] update named_sub_group_size attr test Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/named_sub_group_size.cpp | 15 ++------------- 1 file changed, 2 insertions(+), 13 deletions(-) diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index 81972804a69a4..df96d6ae811be 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -1,5 +1,4 @@ -// 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 +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -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. @@ -15,7 +14,7 @@ struct Functor1 { } }; -// Test attributes get propgated to the kernel. +// Test attribute gets propgated to the kernel. void calls_kernel_1() { // CHECK: FunctionDecl {{.*}}Kernel1 // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic @@ -47,13 +46,3 @@ void calls_kernel_3() { AttrFunc(); }); } - -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(); - }); -} From 847cea427943668b18983119e8e1d9d2335ff796 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 15 Jul 2021 18:16:57 -0700 Subject: [PATCH 14/16] address @elizabeth's review comments Signed-off-by: Soumi Manna --- .../CodeGenSYCL/scheduler-target-fmax-mhz.cpp | 2 +- .../check-direct-attribute-propagation.cpp | 43 +++---- .../intel-max-work-group-size-device.cpp | 4 +- ...eqd-work-group-size-device-direct-prop.cpp | 115 +++++++++--------- .../intel-reqd-work-group-size-device.cpp | 4 +- clang/test/SemaSYCL/named_sub_group_size.cpp | 4 +- ...eqd-work-group-size-device-direct-prop.cpp | 4 - clang/test/SemaSYCL/sycl-esimd-ast.cpp | 2 +- 8 files changed, 87 insertions(+), 91 deletions(-) diff --git a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp index 6bd6c0fbc7cf6..1357347f2208b 100644 --- a/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp +++ b/clang/test/CodeGenSYCL/scheduler-target-fmax-mhz.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64-unknown-unknown-sycldevice -disable-llvm-passes -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 index 95941780fd9bc..53aaea02cda4c 100644 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -90,7 +90,7 @@ int main() { // CHECK: FunctionDecl {{.*}}test_kernel3 // CHECK: SYCLSimdAttr {{.*}} Implicit // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr {{.*}} + // CHECK-NEXT: SYCLSimdAttr // CHECK-NOT: SYCLSimdAttr h.single_task( []() [[intel::sycl_explicit_simd]] { func(); }); @@ -105,7 +105,7 @@ int main() { // CHECK: SYCLIntelNoGlobalWorkOffsetAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 h.single_task( FuncObj1()); @@ -113,7 +113,7 @@ int main() { // CHECK: SYCLIntelNoGlobalWorkOffsetAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 h.single_task( []() [[intel::no_global_work_offset]]{}); @@ -121,7 +121,7 @@ int main() { // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 10 - // CHECK-NEXT: IntegerLiteral{{.*}}10{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 10 h.single_task( FuncObj2()); @@ -129,7 +129,7 @@ int main() { // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 20 - // CHECK-NEXT: IntegerLiteral{{.*}}20{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 20 h.single_task( []() [[intel::scheduler_target_fmax_mhz(20)]]{}); @@ -143,13 +143,13 @@ int main() { // CHECK: SYCLIntelMaxWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 h.single_task( FuncObj3()); @@ -163,13 +163,13 @@ int main() { // CHECK: SYCLIntelMaxWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 h.single_task( []() [[intel::max_work_group_size(8, 8, 8)]]{}); @@ -177,13 +177,13 @@ int main() { // CHECK: ReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral{{.*}}2{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 h.single_task( FuncObj4()); @@ -197,13 +197,13 @@ int main() { // CHECK: ReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 h.single_task( []() [[sycl::reqd_work_group_size(8, 8, 8)]]{}); @@ -211,7 +211,7 @@ int main() { // CHECK: SYCLIntelNumSimdWorkItemsAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 h.single_task( FuncObj5()); @@ -219,7 +219,7 @@ int main() { // CHECK: SYCLIntelNumSimdWorkItemsAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 20 - // CHECK-NEXT: IntegerLiteral{{.*}}20{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 20 h.single_task( []() [[intel::num_simd_work_items(20)]]{}); @@ -249,7 +249,7 @@ int main() { // CHECK: SYCLIntelMaxGlobalWorkDimAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 h.single_task( FuncObj7()); @@ -257,7 +257,7 @@ int main() { // CHECK: SYCLIntelMaxGlobalWorkDimAttr // CHECK-NEXT: ConstantExpr {{.*}} 'int' // CHECK-NEXT: value: Int 0 - // CHECK-NEXT: IntegerLiteral{{.*}}0{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 0 h.single_task( []() [[intel::max_global_work_dim(0)]]{}); @@ -282,6 +282,7 @@ int main() { // CHECK: IntelReqdSubGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 8 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 h.single_task( []() [[sycl::reqd_sub_group_size(8)]]{}); }); 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 02c7d487666dc..906d35c3eaac6 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 -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 +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -triple spir64 -DTRIGGER_ERROR -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -triple spir64 | FileCheck %s #include "sycl.hpp" diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp index ba814f82ef025..8909245dd2b19 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp @@ -68,79 +68,78 @@ class FunctorAttr { int main() { q.submit([&](handler &h) { + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 Functor16 f16; h.single_task(f16); + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 16 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 Functor16x16x16 f16x16x16; h.single_task(f16x16x16); + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 128 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 128 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 128 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 FunctorAttr fattr; h.single_task(fattr); + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 32 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 32 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int -4 + // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 4 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 Functor33 f33; h.single_task(f33); + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 30 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 30 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int -30 + // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 30 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int -30 + // CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' Functor30 f30; h.single_task(f30); }); return 0; } - -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 -// 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{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 128 -// CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name4 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int -4 -// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 30 -// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int -30 -// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' -// CHECK-NEXT: IntegerLiteral{{.*}}30{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int -30 -// CHECK-NEXT: UnaryOperator{{.*}} 'int' prefix '-' #endif // __SYCL_DEVICE_ONLY__ 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 e06ae8bd70d6b..68d299d4d3a4f 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,5 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -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 +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -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 // Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1. diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index df96d6ae811be..bee88cdb5b073 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -14,7 +14,7 @@ struct Functor1 { } }; -// Test attribute gets propgated to the kernel. +// Test attribute gets propagated to the kernel. void calls_kernel_1() { // CHECK: FunctionDecl {{.*}}Kernel1 // CHECK: IntelNamedSubGroupSizeAttr {{.*}} Automatic @@ -36,7 +36,7 @@ void calls_kernel_2() { [[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc -// Test ttribute does not get propgated to the kernel. +// Test attribute does not get propagated to the kernel. void calls_kernel_3() { // CHECK: FunctionDecl {{.*}}Kernel4 // CHECK-NOT: IntelNamedSubGroupSizeAttr {{.*}} diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp index a350f344f00f2..a0c0455aa3449 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp @@ -23,10 +23,6 @@ __attribute__((reqd_work_group_size(4))) void four_yet_again(); // expected-erro [[cl::reqd_work_group_size(4)]] void four_with_more_feeling(); // 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 // TRIGGER_ERROR - - -#ifdef TRIGGER_ERROR class Functor32 { public: // expected-note@+3{{conflicting attribute is here}} diff --git a/clang/test/SemaSYCL/sycl-esimd-ast.cpp b/clang/test/SemaSYCL/sycl-esimd-ast.cpp index 86454f7f06724..a0eae761582dc 100644 --- a/clang/test/SemaSYCL/sycl-esimd-ast.cpp +++ b/clang/test/SemaSYCL/sycl-esimd-ast.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -ast-dump %s | FileCheck %s // Tests for AST of sycl_explicit_simd function attribute in SYCL 1.2.1. From b749df1c8f1299ac9858ea07c95ee433c85be6f1 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 15 Jul 2021 18:27:16 -0700 Subject: [PATCH 15/16] update tests Signed-off-by: Soumi Manna --- .../check-direct-attribute-propagation.cpp | 2 +- .../check-notdirect-attribute-propagation.cpp | 2 +- .../reqd-work-group-size-device-direct-prop.cpp | 16 ++++++++-------- 3 files changed, 10 insertions(+), 10 deletions(-) diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp index 53aaea02cda4c..45fe936643b33 100644 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -Wno-sycl-2017-compat -sycl-std=2020 -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::scheduler_target_fmax_mhz()]], [[intel::num_simd_work_items()]], // [[intel::no_global_work_offset()]], [[intel::max_global_work_dim()]], [[intel::sycl_explicit_simd]], diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp index 03d605080cdd7..bdabd1f40d52d 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 -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 +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2017 -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/reqd-work-group-size-device-direct-prop.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp index a0c0455aa3449..88b8ab6a3a3a1 100644 --- a/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp +++ b/clang/test/SemaSYCL/reqd-work-group-size-device-direct-prop.cpp @@ -46,30 +46,30 @@ class FunctorAttr { int main() { q.submit([&](handler &h) { // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 - // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK: ReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 16 - // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 16 - // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 16 - // CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 16 Functor16x16x16 f16x16x16; h.single_task(f16x16x16); // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 - // CHECK: ReqdWorkGroupSizeAttr {{.*}} + // CHECK: ReqdWorkGroupSizeAttr // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 128 - // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 128 - // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 // CHECK-NEXT: ConstantExpr{{.*}}'int' // CHECK-NEXT: value: Int 128 - // CHECK-NEXT: IntegerLiteral{{.*}}128{{$}} + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 128 FunctorAttr fattr; h.single_task(fattr); }); From ed793a5c24dea7762f5756c1039f653cf5244568 Mon Sep 17 00:00:00 2001 From: Soumi Manna Date: Thu, 15 Jul 2021 18:57:20 -0700 Subject: [PATCH 16/16] update test to use expected-error instead of primary-error Signed-off-by: Soumi Manna --- clang/test/SemaSYCL/named_sub_group_size.cpp | 8 +++----- 1 file changed, 3 insertions(+), 5 deletions(-) diff --git a/clang/test/SemaSYCL/named_sub_group_size.cpp b/clang/test/SemaSYCL/named_sub_group_size.cpp index bee88cdb5b073..d207e73453fa4 100644 --- a/clang/test/SemaSYCL/named_sub_group_size.cpp +++ b/clang/test/SemaSYCL/named_sub_group_size.cpp @@ -1,4 +1,4 @@ -// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -fsyntax-only -ast-dump -verify=expected,primary %s | FileCheck %s +// RUN: %clang_cc1 -internal-isystem %S/Inputs -fsycl-is-device -sycl-std=2020 -fsyntax-only -ast-dump -verify %s | FileCheck %s // Validate the semantic analysis checks for the named_sub_group_size attribute in SYCL 2020 mode. @@ -34,15 +34,13 @@ void calls_kernel_2() { sycl::kernel_single_task(F1); } -[[intel::named_sub_group_size(primary)]] void AttrFunc() {} // #AttrFunc +[[intel::named_sub_group_size(primary)]] void AttrFunc() {} // expected-error{{kernel-called function must have a sub group size that matches the size specified for the kernel}} // Test attribute does not get propagated to the kernel. 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}} + sycl::kernel_single_task([]() { // expected-note{{kernel declared here}} AttrFunc(); }); }