Skip to content
Open
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -6990,7 +6990,10 @@ ExprResult SemaSYCL::ActOnUniqueStableNameExpr(SourceLocation OpLoc,
}

void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) {
if (SYCLKernelFunctions.contains(FD))
// To avoid confusing, we skip issuing warnings for functors that are defined
// but not used.
if (SYCLKernelFunctions.contains(FD) ||
(FD->getOverloadedOperator() == OverloadedOperatorKind::OO_Call))
return;

for (const auto *KernelAttr : std::vector<AttributeCommonInfo *>{
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -101,7 +101,7 @@ class Functor16x2x1 {

class Functor4x4x4 {
public:
[[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
[[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {};
};

// Checking whether propagation of the attribute happens or not, according to the SYCL version.
Expand Down
18 changes: 7 additions & 11 deletions clang/test/SemaSYCL/reqd_work_group_size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -37,16 +37,14 @@ f4() {}
// Catch the easy case where the attributes are all specified at once with
// different arguments.
struct TRIFuncObjGood1 {
// expected-note@+3 {{previous attribute is here}}
// expected-error@+2 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
// expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
// expected-note@+2 {{previous attribute is here}}
// expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
[[sycl::reqd_work_group_size(64)]] [[sycl::reqd_work_group_size(128)]] void operator()() const {}
};

struct TRIFuncObjGood2 {
// expected-note@+3 {{previous attribute is here}}
// expected-error@+2 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
// expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
// expected-note@+2 {{previous attribute is here}}
// expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
[[sycl::reqd_work_group_size(64, 64)]] [[sycl::reqd_work_group_size(128, 128)]] void operator()() const {}
};

Expand All @@ -55,8 +53,7 @@ struct TRIFuncObjGood3 {
operator()() const;
};

[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} \
// expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
void
TRIFuncObjGood3::operator()() const {}

Expand All @@ -77,7 +74,7 @@ class FunctorC {

class Functor32 {
public:
[[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[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 {}
};
Expand Down Expand Up @@ -143,8 +140,7 @@ struct TRIFuncObjBad {
operator()() const;
};

[[sycl::reqd_work_group_size(1, 1, 32)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}} \
// expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[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 {}

Expand Down
22 changes: 8 additions & 14 deletions clang/test/SemaSYCL/sycl-attr-warn-non-kernel.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,7 @@ void Functor16::operator()() const {

class Functor16x16 {
public:
[[sycl::reqd_work_group_size(16,16)]] void operator()() const; // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[sycl::reqd_work_group_size(16,16)]] void operator()() const;
};
void Functor16x16::operator()() const {
}
Expand All @@ -41,21 +41,24 @@ class Functor16x16x16 {

class FunctorSubGroupSize4 {
public:
[[intel::reqd_sub_group_size(4)]] void operator()() const{} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
[[intel::reqd_sub_group_size(4)]] void operator()() const{}
};

class Functor8 {
public:
void operator()() const;
};

[[sycl::reqd_work_group_size(8)]] void Functor8::operator()() const {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[sycl::reqd_work_group_size(8)]] void Functor8::operator()() const {}

int main() {
sycl::queue q;
Functor16x16 f16x16;
FunctorSubGroupSize4 fs4;

// Show that the functors don't issue diagnostics.
Functor8 f8;
FunctorSubGroupSize4 fs4;
Functor16x16 f16x16;
Functor16x16x16 f16x16x16;

q.submit([&](sycl::handler& h) {
h.single_task<class kernel_name>(
Expand All @@ -65,14 +68,5 @@ int main() {
);
});


q.submit([&](sycl::handler &h) {
Functor16 f16;
Functor16x16x16 f16x16x16;
h.single_task<class kernel_name1>(f16); // OK attribute reqd_work_group_size applied to kernel
h.single_task<class kernel_name2>(f16x16x16); // OK attribute reqd_work_group_size applied to kernel
});


return 0;
}
Loading