diff --git a/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp b/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp new file mode 100644 index 0000000000000..171ba0d0c5faf --- /dev/null +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-ast-device.cpp @@ -0,0 +1,201 @@ +// 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; +queue q; + +// Test that checks template parameter support on member function of class template. +template +class KernelFunctor { +public: + [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {} +}; + +void test() { + KernelFunctor<16, 1, 1>(); +} + +// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor +// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition +// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor +// CHECK: ReqdWorkGroupSizeAttr +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 16 +// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 1 +// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + +// Test that checks template parameter support on function. +template +[[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {} + +int check() { + func3<8, 8, 8>(); + return 0; +} +// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3 +// CHECK: FunctionDecl {{.*}} {{.*}} used func3 'void ()' +// CHECK: ReqdWorkGroupSizeAttr +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 8 +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 8 +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} +// CHECK-NEXT: ConstantExpr{{.*}}'int' +// CHECK-NEXT: value: Int 8 +// CHECK: SubstNonTypeTemplateParmExpr {{.*}} +// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} +// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} + +[[sycl::reqd_work_group_size(4)]] void f4x1x1() {} + +class Functor16 { +public: + [[sycl::reqd_work_group_size(16)]] void operator()() const {} +}; + +class Functor16x16x16 { +public: + [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} +}; + +class Functor { +public: + void operator()() const { + f4x1x1(); + } +}; + +class FunctorAttr { +public: + [[sycl::reqd_work_group_size(128, 128, 128)]] void operator()() const {} +}; + +// Test of redeclaration of [[intel::max_work_group_size()]] and [[sycl::reqd_work_group_size()]]. +[[intel::no_global_work_offset]] void func1(); +[[intel::max_work_group_size(4, 4, 4)]] void func1(); +[[sycl::reqd_work_group_size(2, 2, 2)]] void func1() {} + +[[sycl::reqd_work_group_size(32, 32, 32)]] void f32x32x32() {} + +int main() { + q.submit([&](handler &h) { + // 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{{$}} + Functor16 f16; + h.single_task(f16); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name2 + // CHECK: ReqdWorkGroupSizeAttr + // 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{{$}} + Functor f; + h.single_task(f); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name3 + // 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); + + // 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{{$}} + FunctorAttr fattr; + h.single_task(fattr); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name5 + // CHECK: ReqdWorkGroupSizeAttr + // 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{{$}} + h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { + f32x32x32(); + }); + + // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name6 + // CHECK: SYCLIntelNoGlobalWorkOffsetAttr + // CHECK-NEXT: ConstantExpr {{.*}} 'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} + // CHECK: SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited + // 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: 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( + []() { func1(); }); + }); + 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 deleted file mode 100644 index 87d10ac0c0867..0000000000000 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device-direct-prop.cpp +++ /dev/null @@ -1,116 +0,0 @@ -// 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 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: - [[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) { - // 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 64 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 64 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 - // CHECK-NEXT: ConstantExpr{{.*}}'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 - Functor64 f64; - h.single_task(f64); - }); - return 0; -} -#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 1dea68a817583..d93922352fe9b 100644 --- a/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp +++ b/clang/test/SemaSYCL/intel-reqd-work-group-size-device.cpp @@ -1,28 +1,12 @@ -// 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 %s -// Test for AST of reqd_work_group_size kernel attribute in SYCL 1.2.1. +// The test checks support and functionality of reqd_work_group_size kernel attribute in SYCL 2017. #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 [[sycl::reqd_work_group_size(4)]] void f4x1x1() {} // expected-note {{conflicting attribute is here}} // expected-note@-1 {{conflicting attribute is here}} [[sycl::reqd_work_group_size(32)]] void f32x1x1() {} // expected-note {{conflicting attribute is here}} @@ -32,19 +16,7 @@ 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 [[intel::reqd_work_group_size(4, 2, 9)]] void unknown() {} // expected-warning{{unknown attribute 'reqd_work_group_size' ignored}} -#endif // TRIGGER_ERROR - -class Functor16 { -public: - [[sycl::reqd_work_group_size(16)]] void operator()() const {} -}; - -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: @@ -53,117 +25,43 @@ class Functor8 { // expected-error {{conflicting attributes applied to a SYCL ke } }; -class Functor { -public: - void operator()() const { - f4x1x1(); - } -}; +// 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}} -class FunctorAttr { -public: - [[sycl::reqd_work_group_size(128, 128, 128)]] void operator()() const {} -}; +[[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) { - Functor16 f16; - h.single_task(f16); - - Functor f; - h.single_task(f); - - Functor16x16x16 f16x16x16; - h.single_task(f16x16x16); - - FunctorAttr fattr; - h.single_task(fattr); - - h.single_task([]() [[sycl::reqd_work_group_size(32, 32, 32)]] { - f32x32x32(); - }); -#ifdef TRIGGER_ERROR Functor8 f8; - h.single_task(f8); + h.single_task(f8); - h.single_task([]() { // expected-error {{conflicting attributes applied to a SYCL kernel}} + 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}} + 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}} + 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)]][]() { + h.single_task([[sycl::reqd_work_group_size(32, 32, 32)]][]() { f32x32x32(); }); -#endif // TRIGGER_ERROR + h.single_task( + []() { func2(); }); + + h.single_task( + []() { func3(); }); }); 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 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: 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_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' -// 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{{$}} -#endif // __SYCL_DEVICE_ONLY__ diff --git a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp b/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp deleted file mode 100644 index 1a5357a9f0ab2..0000000000000 --- a/clang/test/SemaSYCL/redeclaration-attribute-propagation.cpp +++ /dev/null @@ -1,79 +0,0 @@ -// 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 - -// The test checks redeclaration of [[intel:::max_work_group_size()]] and [[sycl::reqd_work_group_size()]] attributes. -#include "sycl.hpp" - -using namespace cl::sycl; -queue q; - -#ifndef TRIGGER_ERROR -//first case - good case -[[intel::no_global_work_offset]] // expected-no-diagnostics -void -func1(); - -[[intel::max_work_group_size(4, 4, 4)]] void func1(); - -[[sycl::reqd_work_group_size(2, 2, 2)]] void func1() {} - -#else -//second case - 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}} - -//third case - expect error -[[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}} - -// fourth case - expect warning. -[[intel::max_work_group_size(4, 4, 4)]] void func4(); // expected-note {{previous attribute is here}} -[[intel::max_work_group_size(8, 8, 8)]] void func4() {} // expected-warning {{attribute 'max_work_group_size' is already applied with different arguments}} -#endif - -int main() { - q.submit([&](handler &h) { -#ifndef TRIGGER_ERROR - // CHECK-LABEL: FunctionDecl {{.*}} main 'int ()' - // CHECK: `-FunctionDecl {{.*}}test_kernel1 'void ()' - // CHECK: -SYCLIntelNoGlobalWorkOffsetAttr {{.*}} - // CHECK-NEXT: ConstantExpr {{.*}} 'int' - // CHECK-NEXT: value: Int 1 - // CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - // CHECK: -SYCLIntelMaxWorkGroupSizeAttr {{.*}} Inherited - // 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: `-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( - []() { func1(); }); - -#else - h.single_task( - []() { func2(); }); - - h.single_task( - []() { func3(); }); - - h.single_task( - []() { func4(); }); -#endif - }); - return 0; -} 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 bbfa3136cfb2d..7fd80f5dfca6e 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 @@ -1,4 +1,3 @@ -// 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. @@ -8,29 +7,6 @@ using namespace cl::sycl; queue q; -// 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?}} -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 // TRIGGER_ERROR - class Functor16x16x16 { public: [[sycl::reqd_work_group_size(16, 16, 16)]] void operator()() const {} @@ -49,6 +25,16 @@ struct TRIFuncObjGood { [[sycl::reqd_work_group_size(1, 2, 3)]] void TRIFuncObjGood::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 {} +}; + int main() { q.submit([&](handler &h) { // CHECK: FunctionDecl {{.*}} {{.*}}kernel_name1 @@ -91,6 +77,34 @@ int main() { // CHECK-NEXT: value: Int 3 // CHECK-NEXT: IntegerLiteral{{.*}}3{{$}} h.single_task(TRIFuncObjGood()); + + // CHECK: FunctionDecl {{.*}} {{.*}}test_kernel4 + // 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 {{.*}} {{.*}}test_kernel5 + // CHECK: ReqdWorkGroupSizeAttr + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 64 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 64 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 64 + // CHECK-NEXT: ConstantExpr{{.*}}'int' + // CHECK-NEXT: value: Int 1 + // CHECK-NEXT: IntegerLiteral {{.*}} 'int' 1 + Functor64 f64; + h.single_task(f64); }); return 0; } diff --git a/clang/test/SemaSYCL/reqd_work_group_size.cpp b/clang/test/SemaSYCL/reqd_work_group_size.cpp index 318cbf4efeeb0..f4545e62d8d16 100644 --- a/clang/test/SemaSYCL/reqd_work_group_size.cpp +++ b/clang/test/SemaSYCL/reqd_work_group_size.cpp @@ -1,5 +1,7 @@ // RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -verify %s +// The test checks support and functionality of [[sycl::reqd_work_group_size()]] attribute. + // Check the basics. [[sycl::reqd_work_group_size]] void f(); // expected-error {{'reqd_work_group_size' attribute takes at least 1 argument}} [[sycl::reqd_work_group_size(12, 12, 12, 12)]] void f0(); // expected-error {{'reqd_work_group_size' attribute takes no more than 3 arguments}} @@ -69,6 +71,13 @@ class FunctorC { operator()(int) const; }; +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 {} +}; + // Ensure that template arguments behave appropriately based on instantiations. template [[sycl::reqd_work_group_size(N, 1, 1)]] void f6(); // #f6 @@ -128,3 +137,60 @@ struct TRIFuncObjBad { [[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} void TRIFuncObjBad::operator()() const {} + +// Test that checks wrong function template instantiation and ensures that the type +// is checked properly when instantiating from the template definition. + +template +// expected-error@+1 3{{integral constant expression must have integral or unscoped enumeration type, not 'S'}} +[[sycl::reqd_work_group_size(Ty{}, Ty1{}, Ty2{})]] void func() {} + +struct S {}; +void var() { + // expected-note@+1{{in instantiation of function template specialization 'func' requested here}} + func(); +} + +// Test that checks expression is not a constant expression. +// expected-note@+1 3{{declared here}} +int foo(); +// expected-error@+2 3{{expression is not an integral constant expression}} +// expected-note@+1 3{{non-constexpr function 'foo' cannot be used in a constant expression}} +[[sycl::reqd_work_group_size(foo() + 12, foo() + 12, foo() + 12)]] void func1(); + +// Test that checks expression is a constant expression. +constexpr int bar() { return 0; } +[[sycl::reqd_work_group_size(bar() + 12, bar() + 12, bar() + 12)]] void func2(); // OK + +// Test that checks template parameter support on member function of class template. +template +class KernelFunctor { +public: + [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {} +}; + +int main() { + KernelFunctor<16, 1, 1>(); +} +// Test that checks template parameter support on function. +template +[[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {} + +int check() { + func3<8, 8, 8>(); + return 0; +} +// 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?}} + +__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?}} diff --git a/clang/test/SemaSYCL/sycl-device-reqd-work-group-size-template.cpp b/clang/test/SemaSYCL/sycl-device-reqd-work-group-size-template.cpp deleted file mode 100644 index 6e0e5f253bd3c..0000000000000 --- a/clang/test/SemaSYCL/sycl-device-reqd-work-group-size-template.cpp +++ /dev/null @@ -1,86 +0,0 @@ -// RUN: %clang_cc1 -fsycl-is-device -fsyntax-only -ast-dump -verify -pedantic %s | FileCheck %s - -// Test that checks template parameter support for 'sycl::reqd_work_group_size' attribute on sycl device. - -// Test that checks wrong function template instantiation and ensures that the type -// is checked properly when instantiating from the template definition. - -template -// expected-error@+1 3{{integral constant expression must have integral or unscoped enumeration type, not 'S'}} -[[sycl::reqd_work_group_size(Ty{}, Ty1{}, Ty2{})]] void func() {} - -struct S {}; -void var() { - //expected-note@+1{{in instantiation of function template specialization 'func' requested here}} - func(); -} - -// Test that checks expression is not a constant expression. -// expected-note@+1 3{{declared here}} -int foo(); -// expected-error@+2 3{{expression is not an integral constant expression}} -// expected-note@+1 3{{non-constexpr function 'foo' cannot be used in a constant expression}} -[[sycl::reqd_work_group_size(foo() + 12, foo() + 12, foo() + 12)]] void func1(); - -// Test that checks expression is a constant expression. -constexpr int bar() { return 0; } -[[sycl::reqd_work_group_size(bar() + 12, bar() + 12, bar() + 12)]] void func2(); // OK - -// Test that checks template parameter support on member function of class template. -template -class KernelFunctor { -public: - [[sycl::reqd_work_group_size(SIZE, SIZE1, SIZE2)]] void operator()() {} -}; - -int main() { - KernelFunctor<16, 1, 1>(); -} - -// CHECK: ClassTemplateDecl {{.*}} {{.*}} KernelFunctor -// CHECK: ClassTemplateSpecializationDecl {{.*}} {{.*}} class KernelFunctor definition -// CHECK: CXXRecordDecl {{.*}} {{.*}} implicit class KernelFunctor -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 16 -// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}16{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 1 -// CHECK-NEXT: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}1{{$}} - -// Test that checks template parameter support on function. -template -[[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {} - -int check() { - func3<8, 8, 8>(); - return 0; -} - -// CHECK: FunctionTemplateDecl {{.*}} {{.*}} func3 -// CHECK: FunctionDecl {{.*}} {{.*}} used func3 'void ()' -// CHECK: ReqdWorkGroupSizeAttr {{.*}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 8 -// CHECK: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 8 -// CHECK: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}} -// CHECK-NEXT: ConstantExpr{{.*}}'int' -// CHECK-NEXT: value: Int 8 -// CHECK: SubstNonTypeTemplateParmExpr {{.*}} -// CHECK-NEXT: NonTypeTemplateParmDecl {{.*}} -// CHECK-NEXT: IntegerLiteral{{.*}}8{{$}}