Skip to content

Commit

Permalink
[SYCL] Add warning for attributes applied to non-kernel functions (#1…
Browse files Browse the repository at this point in the history
…5154)

According to [5.8.1. Kernel
attributes](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:kernel.attributes),
kernel attributes can only be applied to SYCL kernel functions, but not
to a regular device functions.

Based on this information, this pull request introduces a warning
diagnostic whenever specified attributes are applied to regular device
functions.
  • Loading branch information
DYNIO-INTEL authored Sep 17, 2024
1 parent da480db commit 227af47
Show file tree
Hide file tree
Showing 12 changed files with 156 additions and 25 deletions.
2 changes: 2 additions & 0 deletions clang/include/clang/Basic/DiagnosticSemaKinds.td
Original file line number Diff line number Diff line change
Expand Up @@ -12358,6 +12358,8 @@ def err_bit_cast_type_size_mismatch : Error<
"__builtin_bit_cast source size does not equal destination size (%0 vs %1)">;

// SYCL-specific diagnostics
def warn_sycl_incorrect_use_attribute_non_kernel_function : Warning<
"%0 attribute can only be applied to a SYCL kernel function">, InGroup<SyclStrict>;
def warn_sycl_kernel_num_of_template_params : Warning<
"'sycl_kernel' attribute only applies to a function template with at least"
" two template parameters">, InGroup<IgnoredAttributes>;
Expand Down
8 changes: 8 additions & 0 deletions clang/include/clang/Sema/SemaSYCL.h
Original file line number Diff line number Diff line change
Expand Up @@ -259,6 +259,8 @@ class SemaSYCL : public SemaBase {
// useful notes that shows where the kernel was called.
bool DiagnosingSYCLKernel = false;

llvm::DenseSet<const FunctionDecl *> SYCLKernelFunctions;

public:
SemaSYCL(Sema &S);

Expand Down Expand Up @@ -300,6 +302,10 @@ class SemaSYCL : public SemaBase {
void addSyclDeviceDecl(Decl *d) { SyclDeviceDecls.insert(d); }
llvm::SetVector<Decl *> &syclDeviceDecls() { return SyclDeviceDecls; }

void addSYCLKernelFunction(const FunctionDecl *FD) {
SYCLKernelFunctions.insert(FD);
}

/// Lazily creates and returns SYCL integration header instance.
SYCLIntegrationHeader &getSyclIntegrationHeader() {
if (SyclIntHeader == nullptr)
Expand Down Expand Up @@ -375,6 +381,8 @@ class SemaSYCL : public SemaBase {
SourceLocation Loc,
DeviceDiagnosticReason Reason);

void performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD);

/// Tells whether given variable is a SYCL explicit SIMD extension's "private
/// global" variable - global variable in the private address space.
bool isSYCLEsimdPrivateGlobal(VarDecl *VDecl) {
Expand Down
4 changes: 4 additions & 0 deletions clang/lib/Sema/Sema.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -1909,6 +1909,10 @@ class DeferredDiagnosticsEmitter
void checkFunc(SourceLocation Loc, FunctionDecl *FD) {
auto &Done = DoneMap[InOMPDeviceContext > 0 ? 1 : 0];
FunctionDecl *Caller = UsePath.empty() ? nullptr : UsePath.back();

if (!Caller && S.LangOpts.SYCLIsDevice)
S.SYCL().performSYCLDelayedAttributesAnalaysis(FD);

if ((!ShouldEmitRootNode && !S.getLangOpts().OpenMP && !Caller) ||
S.shouldIgnoreInHostDeviceCheck(FD) || InUsePath.count(FD))
return;
Expand Down
18 changes: 18 additions & 0 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -882,6 +882,8 @@ class SingleDeviceFunctionTracker {
// having a kernel lambda with a lambda call inside of it.
KernelBody = CurrentDecl;
}
if (KernelBody)
Parent.SemaSYCLRef.addSYCLKernelFunction(KernelBody);
}

// Recurse.
Expand Down Expand Up @@ -6852,3 +6854,19 @@ ExprResult SemaSYCL::ActOnUniqueStableNameExpr(SourceLocation OpLoc,

return BuildUniqueStableNameExpr(OpLoc, LParen, RParen, TSI);
}

void SemaSYCL::performSYCLDelayedAttributesAnalaysis(const FunctionDecl *FD) {
if (SYCLKernelFunctions.contains(FD))
return;

for (const auto *KernelAttr : std::vector<AttributeCommonInfo *>{
FD->getAttr<SYCLReqdWorkGroupSizeAttr>(),
FD->getAttr<IntelReqdSubGroupSizeAttr>(),
FD->getAttr<SYCLWorkGroupSizeHintAttr>(),
FD->getAttr<VecTypeHintAttr>()}) {
if (KernelAttr)
Diag(KernelAttr->getLoc(),
diag::warn_sycl_incorrect_use_attribute_non_kernel_function)
<< KernelAttr;
}
}
10 changes: 7 additions & 3 deletions clang/test/SemaSYCL/check-work-group-size-hint-device.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,8 @@

// Produce a conflicting attribute warning when the args are different.
[[sycl::work_group_size_hint(4, 1, 1)]] void f3(); // expected-note {{previous attribute is here}}
[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}}
[[sycl::work_group_size_hint(1, 1, 32)]] void f3() {} // expected-warning {{attribute 'work_group_size_hint' is already applied with different arguments}} \
// expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}

// 1 and 2 dim versions
[[sycl::work_group_size_hint(2)]] void f4(); // ok
Expand Down Expand Up @@ -70,10 +71,13 @@ void instantiate() {
f8<0>(); // expected-note {{in instantiation}}
#endif

// expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
f9<1, 1, 1>(); // OK, args are the same on the redecl.

// expected-warning@#f9 {{attribute 'work_group_size_hint' is already applied with different arguments}}
// expected-note@#f9prev {{previous attribute is here}}
// expected-warning@#f9prev {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}

f9<1, 2, 3>(); // expected-note {{in instantiation}}
}

Expand All @@ -97,14 +101,14 @@ class Functor16x2x1 {

class Functor4x4x4 {
public:
[[sycl::work_group_size_hint(4, 4, 4)]] void operator()() const {};
[[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}}
};

// Checking whether propagation of the attribute happens or not, according to the SYCL version.
#if defined(EXPECT_PROP) // if attribute is propagated, then we expect errors here
void f8x8x8(){};
#else // otherwise no error
[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){};
[[sycl::work_group_size_hint(8, 8, 8)]] void f8x8x8(){}; // expected-warning {{'work_group_size_hint' attribute can only be applied to a SYCL kernel function}}
#endif
class FunctorNoProp {
public:
Expand Down
1 change: 0 additions & 1 deletion clang/test/SemaSYCL/device_has.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,6 @@ enum class aspect {

[[sycl::device_has("123")]] void func1() {} // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}}
[[sycl::device_has(fake_cl::sycl::aspect::aspect1)]] void func2() {} // expected-error{{'device_has' attribute argument is invalid; argument must be device aspect of type sycl::aspect}}

[[sycl::device_has(sycl::aspect::cpu)]] void func3(); // expected-note{{previous attribute is here}}
[[sycl::device_has(sycl::aspect::gpu)]] void func3() {} // expected-warning{{attribute 'device_has' is already applied}}

Expand Down
8 changes: 4 additions & 4 deletions clang/test/SemaSYCL/intel-max-work-group-size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -70,15 +70,15 @@ void instantiate() {
// a declaration along with [[sycl::reqd_work_group_size(X1, Y1, Z1)]]
// attribute, check to see if values of reqd_work_group_size arguments are
// equal or less than values coming from max_work_group_size attribute.
[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}}
[[sycl::reqd_work_group_size(64, 64, 64)]] // expected-note {{conflicting attribute is here}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[intel::max_work_group_size(64, 16, 64)]] // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
void
f9() {}

[[intel::max_work_group_size(4, 4, 4)]] void f10();
[[sycl::reqd_work_group_size(2, 2, 2)]] void f10(); // OK

[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // OK
[[sycl::reqd_work_group_size(2, 2, 2)]] [[intel::max_work_group_size(4, 4, 4)]] void f11() {} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}

[[sycl::reqd_work_group_size(64, 64, 64)]] void f12(); // expected-note {{conflicting attribute is here}}
[[intel::max_work_group_size(16, 16, 16)]] void f12(); // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
Expand All @@ -91,14 +91,14 @@ f13() {}
[[sycl::reqd_work_group_size(64, 64, 64)]] void f14(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}

[[cl::reqd_work_group_size(1, 2, 3)]] // expected-warning {{attribute 'cl::reqd_work_group_size' is deprecated}} \
// expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}}
// expected-note {{did you mean to use 'sycl::reqd_work_group_size' instead?}} // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[intel::max_work_group_size(1, 2, 3)]] void
f15() {} // OK

[[intel::max_work_group_size(2, 3, 7)]] void f16(); // expected-note {{conflicting attribute is here}}
[[sycl::reqd_work_group_size(7, 3, 2)]] void f16(); // expected-error{{'reqd_work_group_size' attribute conflicts with 'max_work_group_size' attribute}}

[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // OK
[[intel::max_work_group_size(1, 2, 3)]] [[sycl::reqd_work_group_size(1, 2, 3)]] void f17(){}; // expected-warning {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}

[[sycl::reqd_work_group_size(16)]] // expected-note {{conflicting attribute is here}}
[[intel::max_work_group_size(16, 1, 1)]] void // expected-error {{'max_work_group_size' attribute conflicts with 'reqd_work_group_size' attribute}}
Expand Down
18 changes: 11 additions & 7 deletions clang/test/SemaSYCL/reqd-sub-group-size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -33,17 +33,19 @@ int main() {
return 0;
}
[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B();
[[intel::reqd_sub_group_size(16)]] void A() {
[[intel::reqd_sub_group_size(16)]] void A() // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
{
}

[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() {
[[intel::reqd_sub_group_size(16)]] SYCL_EXTERNAL void B() { // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
A();
}
// expected-note@+1 {{conflicting attribute is here}}
[[intel::reqd_sub_group_size(2)]] void sg_size2() {}
[[intel::reqd_sub_group_size(2)]] void sg_size2() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}

// expected-note@+2 {{conflicting attribute is here}}
// expected-error@+1 {{conflicting attributes applied to a SYCL kernel}}
// expected-note@+3 {{conflicting attribute is here}}
// expected-error@+2 {{conflicting attributes applied to a SYCL kernel}}
// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
[[intel::reqd_sub_group_size(4)]] __attribute__((sycl_device)) void sg_size4() {
sg_size2();
}
Expand All @@ -67,7 +69,7 @@ int main() {

// No diagnostic is emitted because the arguments match.
[[intel::reqd_sub_group_size(12)]] void same();
[[intel::reqd_sub_group_size(12)]] void same() {} // OK
[[intel::reqd_sub_group_size(12)]] void same() {} // expected-warning {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}

// No diagnostic because the attributes are synonyms with identical behavior.
[[sycl::reqd_sub_group_size(12)]] void same(); // OK
Expand Down Expand Up @@ -117,10 +119,12 @@ int check() {

// Test that checks template parameter support on function.
template <int N>
// expected-error@+1{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
// expected-error@+2{{'reqd_sub_group_size' attribute requires a positive integral compile time constant expression}}
// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
[[intel::reqd_sub_group_size(N)]] void func3() {}

template <int N>
// expected-warning@+1 {{'reqd_sub_group_size' attribute can only be applied to a SYCL kernel function}}
[[intel::reqd_sub_group_size(4)]] void func4(); // expected-note {{previous attribute is here}}

template <int N>
Expand Down
30 changes: 20 additions & 10 deletions clang/test/SemaSYCL/reqd_work_group_size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -22,28 +22,31 @@ class Functor30 {

// Tests for 'reqd_work_group_size' attribute duplication.
// No diagnostic is emitted because the arguments match. Duplicate attribute is silently ignored.
// expected-warning@+1 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
[[sycl::reqd_work_group_size(6, 6, 6)]] [[sycl::reqd_work_group_size(6, 6, 6)]] void f2() {}

// No diagnostic is emitted because the arguments match.
[[sycl::reqd_work_group_size(32, 32, 32)]] void f3();
[[sycl::reqd_work_group_size(32, 32, 32)]] void f3(); // OK

// Produce a conflicting attribute warning when the args are different.
[[sycl::reqd_work_group_size(6, 6, 6)]] // expected-note {{previous attribute is here}}
[[sycl::reqd_work_group_size(6, 6, 6)]] // 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(16, 16, 16)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
f4() {}

// Catch the easy case where the attributes are all specified at once with
// different arguments.
struct TRIFuncObjGood1 {
// expected-note@+2 {{previous attribute is here}}
// expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
// 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}}
[[sycl::reqd_work_group_size(64)]] [[sycl::reqd_work_group_size(128)]] void operator()() const {}
};

struct TRIFuncObjGood2 {
// expected-note@+2 {{previous attribute is here}}
// expected-error@+1 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
// 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}}
[[sycl::reqd_work_group_size(64, 64)]] [[sycl::reqd_work_group_size(128, 128)]] void operator()() const {}
};

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

[[sycl::reqd_work_group_size(4, 4)]] // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
[[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}}
void
TRIFuncObjGood3::operator()() const {}

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

class Functor32 {
public:
[[sycl::reqd_work_group_size(32, 1, 1)]] // expected-note {{previous attribute is here}}
[[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(1, 1, 32)]] void // expected-error{{attribute 'reqd_work_group_size' is already applied with different arguments}}
operator()() const {}
};
Expand Down Expand Up @@ -105,16 +109,18 @@ void instantiate() {
f7<1, 1, 1>(); // OK, args are the same on the redecl.
// expected-error@#f7 {{attribute 'reqd_work_group_size' is already applied with different arguments}}
// expected-note@#f7prev {{previous attribute is here}}
// expected-warning@#f7prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
// expected-warning@#f7prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
f7<2, 2, 2>(); // expected-note {{in instantiation}}
}

// Tests for 'reqd_work_group_size' attribute duplication.

[[sycl::reqd_work_group_size(8)]] // expected-note {{previous attribute is here}}
[[sycl::reqd_work_group_size(8)]] // 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(1, 1, 8)]] void // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}
f8(){};

[[sycl::reqd_work_group_size(32, 32, 1)]] // expected-note {{previous attribute is here}}
[[sycl::reqd_work_group_size(32, 32, 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, 32)]] void f9() {} // expected-error {{attribute 'reqd_work_group_size' is already applied with different arguments}}

// Test that template redeclarations also get diagnosed properly.
Expand All @@ -127,6 +133,8 @@ void test() {
f10<64, 1, 1>(); // OK, args are the same on the redecl.
// expected-error@#f10err {{attribute 'reqd_work_group_size' is already applied with different arguments}}
// expected-note@#f10prev {{previous attribute is here}}
// expected-warning@#f10prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
// expected-warning@#f10prev {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
f10<1, 1, 64>(); // expected-note {{in instantiation}}
}

Expand All @@ -135,7 +143,8 @@ 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}}
[[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}}
void
TRIFuncObjBad::operator()() const {}

Expand Down Expand Up @@ -174,6 +183,7 @@ int main() {
KernelFunctor<16, 1, 1>();
}
// Test that checks template parameter support on function.
// expected-warning@+2 {{'reqd_work_group_size' attribute can only be applied to a SYCL kernel function}}
template <int N, int N1, int N2>
[[sycl::reqd_work_group_size(N, N1, N2)]] void func3() {}

Expand Down
2 changes: 2 additions & 0 deletions clang/test/SemaSYCL/sub-group-size.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -123,9 +123,11 @@ void calls_kernel_4() {
sycl::kernel_single_task<class Kernel4>([]() { // #Kernel4
// integer-error@#AttrFunc2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
// integer-note@#Kernel4{{kernel declared here}}
// expected-warning@#AttrFunc2 {{'sub_group_size' attribute can only be applied to a SYCL kernel function}}
AttrFunc2();
// integer-error@#AttrExternalDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
// integer-note@#Kernel4{{kernel declared here}}
// expected-warning@#AttrExternalDefined2 {{'sub_group_size' attribute can only be applied to a SYCL kernel function}}
AttrExternalDefined2();
// integer-error@#AttrExternalNotDefined2{{kernel-called function must have a sub group size that matches the size specified for the kernel}}
// integer-note@#Kernel4{{kernel declared here}}
Expand Down
Loading

0 comments on commit 227af47

Please sign in to comment.