From 83148d028f40ac1d8024600b03b09d128c1afd76 Mon Sep 17 00:00:00 2001 From: Andrei Zibrov Date: Fri, 21 Jun 2024 05:57:51 -0700 Subject: [PATCH] [SYCL] rest of -Wno-sycl-2017-compat removal in SemaSYCL tests --- .../check-direct-attribute-propagation.cpp | 293 ------------------ .../check-notdirect-attribute-propagation.cpp | 54 ---- .../check-work-group-size-hint-device.cpp | 28 +- .../intel-reqd-work-group-size-device.cpp | 67 ---- clang/test/SemaSYCL/intel-restrict.cpp | 10 +- clang/test/SemaSYCL/reqd-sub-group-size.cpp | 42 +-- .../SemaSYCL/reqd-work-group-size-device.cpp | 186 ----------- .../test/SemaSYCL/sycl-2017-future-compat.cpp | 13 - clang/test/SemaSYCL/sycl-esimd.cpp | 17 +- 9 files changed, 28 insertions(+), 682 deletions(-) delete mode 100644 clang/test/SemaSYCL/check-direct-attribute-propagation.cpp delete mode 100644 clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp delete mode 100644 clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp delete mode 100644 clang/test/SemaSYCL/reqd-work-group-size-device.cpp delete mode 100644 clang/test/SemaSYCL/sycl-2017-future-compat.cpp diff --git a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp b/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp deleted file mode 100644 index f40bf981ea606..0000000000000 --- a/clang/test/SemaSYCL/check-direct-attribute-propagation.cpp +++ /dev/null @@ -1,293 +0,0 @@ -// 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]], -// [[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" - -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() {} - -[[sycl::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() {} - -[[sycl::reqd_sub_group_size(4)]] void func8() {} - -class Functor { -public: - void operator()() const { - func8(); - } -}; - -class Functor1 { -public: - [[sycl::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: AsmLabelAttr {{.*}} Implicit - // CHECK-NEXT: SYCLSimdAttr - h.single_task( - FuncObj()); - // CHECK: FunctionDecl {{.*}}test_kernel2 - // CHECK: SYCLSimdAttr {{.*}} Implicit - // CHECK-NEXT: SYCLKernelAttr {{.*}} Implicit - // CHECK-NEXT: AsmLabelAttr {{.*}} 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: AsmLabelAttr {{.*}} 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 {{.*}} 'int' 1 - h.single_task( - FuncObj1()); - - // CHECK: FunctionDecl {{.*}}test_kernel6 - // CHECK: SYCLIntelNoGlobalWorkOffsetAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 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 {{.*}} 'int' 10 - h.single_task( - FuncObj2()); - - // CHECK: FunctionDecl {{.*}}test_kernel8 - // CHECK: SYCLIntelSchedulerTargetFmaxMhzAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 20 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 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 {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 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 {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - h.single_task( - []() [[intel::max_work_group_size(8, 8, 8)]]{}); - - // CHECK: FunctionDecl {{.*}}test_kernel13 - // CHECK: SYCLReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 2 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 2 - h.single_task( - FuncObj4()); - - // Test attribute is not propagated. - // CHECK: FunctionDecl {{.*}}test_kernel14 - // CHECK-NOT: SYCLReqdWorkGroupSizeAttr - h.single_task( - []() { func4(); }); - - // CHECK: FunctionDecl {{.*}}test_kernel15 - // CHECK: SYCLReqdWorkGroupSizeAttr - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 8 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 8 - h.single_task( - []() [[sycl::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 {{.*}} 'int' 8 - h.single_task( - FuncObj5()); - - // CHECK: FunctionDecl {{.*}}test_kernel17 - // CHECK: SYCLIntelNumSimdWorkItemsAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 20 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 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 {{.*}} 'int' 1 - h.single_task( - FuncObj7()); - - // CHECK: FunctionDecl {{.*}}test_kernel23 - // CHECK: SYCLIntelMaxGlobalWorkDimAttr - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 0 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 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); - - // CHECK: FunctionDecl {{.*}}test_kernel27 - // 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)]]{}); - }); - return 0; -} diff --git a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp b/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp deleted file mode 100644 index bdabd1f40d52d..0000000000000 --- a/clang/test/SemaSYCL/check-notdirect-attribute-propagation.cpp +++ /dev/null @@ -1,54 +0,0 @@ -// 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 -triple spir64 | FileCheck %s - -#ifndef TRIGGER_ERROR -[[intel::no_global_work_offset]] void not_direct_one() {} // expected-no-diagnostics - -[[intel::reqd_sub_group_size(1)]] void func_one() { - not_direct_one(); -} - -#else -[[sycl::reqd_work_group_size(2, 2, 2)]] void not_direct_two() {} // expected-note 2 {{conflicting attribute is here}} - -[[intel::max_work_group_size(1, 1, 1)]] // expected-note 3 {{conflicting attribute is here}} -void -func_two() { - not_direct_two(); -} - -[[sycl::reqd_work_group_size(4, 4, 4)]] // expected-note 1 {{conflicting attribute is here}} -void -func_three() { - not_direct_two(); -} -#endif - -template -[[clang::sycl_kernel]] void __my_kernel__(const Type &bar) { - bar(); -#ifndef TRIGGER_ERROR - func_one(); -#else - func_two(); - func_three(); -#endif -} - -template -void parallel_for(Type lambda) { - __my_kernel__(lambda); -} - -void invoke_foo2() { -#ifndef TRIGGER_ERROR - // CHECK-LABEL: FunctionDecl {{.*}} invoke_foo2 'void ()' - // CHECK: `-FunctionDecl {{.*}}KernelName 'void ()' - // CHECK: -IntelReqdSubGroupSizeAttr {{.*}} - // CHECK: `-SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - parallel_for([]() {}); -#else - parallel_for([]() {}); // expected-error 3 {{conflicting attributes applied to a SYCL kernel or SYCL_EXTERNAL function}} -#endif -} diff --git a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp index b989bb3d8bb25..4754cd229e819 100644 --- a/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp +++ b/clang/test/SemaSYCL/check-work-group-size-hint-device.cpp @@ -1,10 +1,10 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -DEXPECT_PROP -DTRIGGER_ERROR %s -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -ast-dump -DEXPECT_PROP %s | FileCheck %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify -DEXPECT_PROP -DTRIGGER_ERROR %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2020 -ast-dump -DEXPECT_PROP %s | FileCheck %s // 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 work_group_size_hint kernel attribute in SYCL 1.2.1. and SYCL 2020 modes. +// Test for AST of work_group_size_hint kernel attribute in SYCL 2020 modes. #include "sycl.hpp" // Check the basics. @@ -133,28 +133,6 @@ void invoke() { // CHECK-NEXT: value: Int 1 // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // Checking that attributes are propagated to the kernel from functions in SYCL 1.2.1 mode. -#ifdef EXPECT_PROP - h.single_task([=]() { - f4x4x4(); - }); -#else - // Otherwise using a functor that has the required attributes - h.single_task(f4x4x4); -#endif - // CHECK: FunctionDecl {{.*}} {{.*}}kernel_2 - // CHECK: SYCLWorkGroupSizeHintAttr {{.*}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 4 - // CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} - - // Check that conflicts are reported if the attribute is propagated in SYCL 1.2.1 mode. FunctorNoProp fNoProp; h.single_task(fNoProp); diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp deleted file mode 100644 index 88ac8380a69b7..0000000000000 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ /dev/null @@ -1,67 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -sycl-std=2017 -Wno-sycl-2017-compat -fsyntax-only -verify %s - -// The test checks support and functionality of reqd_work_group_size kernel attribute in SYCL 2017. - -#include "sycl.hpp" - -using namespace sycl; -queue q; - -[[sycl::reqd_work_group_size(4)]] void f4() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32)]] void f32() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16)]] void f16() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16, 16)]] void f16x16() {} // expected-note {{conflicting attribute is here}} - -[[sycl::reqd_work_group_size(32, 32)]] void f32x32() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} - -[[intel::reqd_work_group_size(4, 2, 9)]] void unknown() {} // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} - -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[sycl::reqd_work_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - f4(); - } -}; - -// Tests of redeclaration of [[intel::max_work_group_size()]] and [[sycl::reqd_work_group_size()]] - expect error -[[intel::max_work_group_size(4, 4, 4)]] void func2(); // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(8, 8, 8)]] void func2() {} // expected-error {{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}} - -[[sycl::reqd_work_group_size(4, 4, 4)]] void func3(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 1)]] void func3() {} // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} - -int main() { - q.submit([&](handler &h) { - Functor8 f8; - h.single_task(f8); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f4(); - f32(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f16(); - f16x16(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f32x32x32(); - f32x32(); - }); - - // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { - f32x32x32(); - }); - - h.single_task( - []() { func2(); }); - - h.single_task( - []() { func3(); }); - }); - return 0; -} diff --git a/clang/test/SemaSYCL/intel-restrict.cpp b/clang/test/SemaSYCL/intel-restrict.cpp index f8e7a670fddd6..b698f464d8822 100644 --- a/clang/test/SemaSYCL/intel-restrict.cpp +++ b/clang/test/SemaSYCL/intel-restrict.cpp @@ -1,7 +1,5 @@ -// 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() {} +// RUN: %clang_cc1 %s -fsyntax-only -fsycl-is-device -sycl-std=2020 -triple spir64 -DCHECKDIAG -verify +// RUN: %clang_cc1 %s -fsyntax-only -ast-dump -fsycl-is-device -sycl-std=2020 -triple spir64 | FileCheck %s struct FuncObj { [[intel::kernel_args_restrict]] void operator()() const {} @@ -26,8 +24,4 @@ int main() { kernel( []() [[intel::kernel_args_restrict]] {}); - // CHECK-LABEL: FunctionDecl {{.*}}test_kernel3 - // CHECK: SYCLIntelKernelArgsRestrictAttr - kernel( - []() { func_do_not_ignore(); }); } diff --git a/clang/test/SemaSYCL/reqd-sub-group-size.cpp b/clang/test/SemaSYCL/reqd-sub-group-size.cpp index 85afcf97388ff..23d2315f67dc0 100644 --- a/clang/test/SemaSYCL/reqd-sub-group-size.cpp +++ b/clang/test/SemaSYCL/reqd-sub-group-size.cpp @@ -1,32 +1,34 @@ -// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify -pedantic %s +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -fsyntax-only -sycl-std=2020 -verify -pedantic %s // The test checks functionality of [[intel::reqd_sub_group_size()]] attribute on SYCL kernel. -#include "sycl.hpp" -using namespace sycl; -queue q; +#include "sycl.hpp" //clang/test/SemaSYCL/Inputs/sycl.hpp -[[intel::reqd_sub_group_size(4)]] void foo() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[intel::reqd_sub_group_size(32)]] void baz() {} // expected-note {{conflicting attribute is here}} +sycl::queue q; -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[intel::reqd_sub_group_size(8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - foo(); - } +// Kernel defined as a named function object +class KernelFunctor1 { + public: + [[sycl::reqd_work_group_size(16)]] void operator()() const {}; +}; + +// Kernel defined as a named function object +class KernelFunctor2 { + public: + void operator() [[sycl::reqd_work_group_size(16)]] () const {}; }; int main() { - q.submit([&](handler &h) { - Functor8 f8; - h.single_task(f8); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - foo(); - baz(); - }); + // Kernel defined as a lambda + q.submit([&](sycl::handler& h) { + KernelFunctor1 kf1; + KernelFunctor2 kf2; + h.single_task(kf1); + h.single_task(kf2); + h.single_task( + []()[[sycl::reqd_work_group_size(16)]]{} + ); }); return 0; } diff --git a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp b/clang/test/SemaSYCL/reqd-work-group-size-device.cpp deleted file mode 100644 index f65a1b59909c4..0000000000000 --- a/clang/test/SemaSYCL/reqd-work-group-size-device.cpp +++ /dev/null @@ -1,186 +0,0 @@ -// 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 sycl; -queue q; - -[[sycl::reqd_work_group_size(4, 1, 1)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} -// expected-note@-1 {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32, 1, 1)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} - -[[sycl::reqd_work_group_size(16, 1, 1)]] void f16x1x1() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(16, 16, 1)]] void f16x16x1() {} // expected-note {{conflicting attribute is here}} - -[[sycl::reqd_work_group_size(32, 32, 1)]] void f32x32x1() {} // expected-note {{conflicting attribute is here}} -[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} // expected-note {{conflicting attribute is here}} - -// No diagnostic because the attributes are synonyms with identical behavior. -[[sycl::reqd_work_group_size(4, 4, 4)]] void four1(); -[[sycl::reqd_work_group_size(4, 4, 4)]] void four1(); // OK -[[sycl::reqd_work_group_size(4, 4)]] void four2(); -[[sycl::reqd_work_group_size(4, 4)]] void four2(); // OK -[[sycl::reqd_work_group_size(4)]] void four3(); -[[sycl::reqd_work_group_size(4)]] void four3(); // OK - -#ifdef TRIGGER_ERROR -// Althrough optional values are effectively representing 1's, the attributes -// should be considered different when there is a different in the arguments -// given. -[[sycl::reqd_work_group_size(4)]] void four4(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(4, 1)]] void four4(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -[[sycl::reqd_work_group_size(4)]] void four5(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(4, 1, 1)]] void four5(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -[[sycl::reqd_work_group_size(4)]] void four6(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 4)]] void four6(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -[[sycl::reqd_work_group_size(4)]] void four7(); // expected-note {{previous attribute is here}} -[[sycl::reqd_work_group_size(1, 1, 4)]] void four7(); // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} -#endif - -// 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 - -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: - [[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} - [[sycl::reqd_work_group_size(1, 1, 32)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} - operator()() const {} -}; -#endif -class Functor16x16x16 { -public: - [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} -}; - -class Functor8 { // expected-error {{conflicting attributes applied to a SYCL kernel}} -public: - [[sycl::reqd_work_group_size(1, 1, 8)]] void operator()() const { // expected-note {{conflicting attribute is here}} - f4x1x1(); - } -}; - -class Functor { -public: - void operator()() const { - f4x1x1(); - } -}; - -int main() { - q.submit([&](handler &h) { - Functor16 f16; - h.single_task(f16); - - Functor f; - h.single_task(f); - - Functor16x16x16 f16x16x16; - h.single_task(f16x16x16); - - h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32), sycl::reqd_work_group_size(32, 32, 32)]] { - f32x32x32(); - }); - -#ifdef TRIGGER_ERROR - Functor8 f8; - h.single_task(f8); - - Functor32 f32; - h.single_task(f32); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f4x1x1(); - f32x1x1(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f16x1x1(); - f16x16x1(); - }); - - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} - f32x32x32(); - f32x32x1(); - }); - - // expected-error@+1 {{expected variable name or 'this' in lambda capture list}} - h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { - f32x32x32(); - }); - -#endif - // Ignore duplicate attribute. - h.single_task( - []() [[sycl::reqd_work_group_size(2, 2, 2), - sycl::reqd_work_group_size(2, 2, 2)]] {}); - }); - return 0; -} - -// CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// 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: SYCLReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 4 -// CHECK-NEXT: IntegerLiteral{{.*}}4{{$}} -// 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_name3 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// 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_name5 -// CHECK: SYCLReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 32 -// CHECK-NEXT: IntegerLiteral{{.*}}32{{$}} -// -// CHECK: FunctionDecl {{.*}}test_kernel11 -// CHECK: SYCLReqdWorkGroupSizeAttr -// 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{{$}} -// CHECK-NOT: SYCLReqdWorkGroupSizeAttr diff --git a/clang/test/SemaSYCL/sycl-2017-future-compat.cpp b/clang/test/SemaSYCL/sycl-2017-future-compat.cpp deleted file mode 100644 index 477c01367fc8f..0000000000000 --- a/clang/test/SemaSYCL/sycl-2017-future-compat.cpp +++ /dev/null @@ -1,13 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -verify %s -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2017 -Wno-sycl-2017-compat -verify=sycl-2017 %s - -// We do not expect any diagnostics from this file when disabling future compat -// warnings. -// sycl-2017-no-diagnostics - -// Test that we get compatibility warnings when using a SYCL 2020 attribute -// spelling while not in SYCL 2020 mode. -[[sycl::reqd_work_group_size(1, 1, 1)]] void f1(); // expected-warning {{use of attribute 'reqd_work_group_size' is a SYCL 2020 extension}} -[[sycl::work_group_size_hint(1, 1, 1)]] void f2(); // expected-warning {{use of attribute 'work_group_size_hint' is a SYCL 2020 extension}} -[[sycl::reqd_sub_group_size(1)]] void f3(); // expected-warning {{use of attribute 'reqd_sub_group_size' is a SYCL 2020 extension}} -[[sycl::vec_type_hint(int)]] void f4(); // expected-warning {{use of attribute 'vec_type_hint' is a SYCL 2020 extension}} diff --git a/clang/test/SemaSYCL/sycl-esimd.cpp b/clang/test/SemaSYCL/sycl-esimd.cpp index 5153aa1ea3b48..f39964107e948 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 -sycl-std=2017 -Wno-sycl-2017-compat -verify %s +// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -sycl-std=2020 -verify %s // This test checks specifics of semantic analysis of ESIMD kernels. @@ -19,14 +19,6 @@ void kernel0(const F &f) __attribute__((sycl_kernel)) { f(); } -// expected-note@+1{{conflicting attribute is here}} -[[intel::reqd_sub_group_size(2)]] void g0() {} - -void test0() { - // expected-error@+2{{conflicting attributes applied to a SYCL kernel}} - // expected-note@+1{{conflicting attribute is here}} - kernel0([=]() __attribute__((sycl_explicit_simd)) { g0(); }); -} // -- Usual kernel can't call ESIMD function template @@ -34,13 +26,6 @@ void kernel1(const F &f) __attribute__((sycl_kernel)) { f(); } -// expected-note@+1{{attribute is here}} -__attribute__((sycl_explicit_simd)) void g1() {} - -void test1() { - // expected-error@+1{{SYCL kernel without 'sycl_explicit_simd' attribute can't call a function with this attribute}} - kernel1([=]() { g1(); }); -} // ----------- Positive tests