Skip to content

[sycl] [clang] Add sycl global var attribute #3746

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 17 commits into from
Jun 9, 2021
Merged

[sycl] [clang] Add sycl global var attribute #3746

Show file tree
Hide file tree
Changes from 9 commits
Commits
Show all changes
17 commits
Select commit Hold shift + click to select a range
682277f
[clang] Add sycl global var attribute
jtmott-intel May 13, 2021
2c6f8fa
Update clang/lib/Sema/SemaExpr.cpp
jtmott-intel May 13, 2021
4184aa4
Use SimpleHandler = 1 and revert changes to SemaDeclAttr
jtmott-intel May 13, 2021
3f7887f
Wrapped attribute in pre-defined macro
jtmott-intel May 13, 2021
63c4f17
Add tests and docs, and remove predefined macro
jtmott-intel May 15, 2021
8b721d9
Update clang/include/clang/Basic/AttrDocs.td
jtmott-intel May 17, 2021
08d0c67
Update clang/include/clang/Basic/AttrDocs.td
jtmott-intel May 17, 2021
5c0f8af
Expand documentation and tests.
jtmott-intel May 18, 2021
a1ac1c5
Update clang/include/clang/Basic/AttrDocs.td
jtmott-intel May 18, 2021
a046f39
Add anonymous union to tests
jtmott-intel May 19, 2021
6bb98b9
Disallowed local vars, and upgraded diagnostic to an error
jtmott-intel May 20, 2021
47ce16a
Restrict attribute to system headers, or it has no effect
jtmott-intel May 27, 2021
a349032
Changed to non-simple handler, added templates to test
jtmott-intel Jun 2, 2021
79d4ed7
clang format
jtmott-intel Jun 2, 2021
e9b7b2a
Simulated system include, and added macro test
jtmott-intel Jun 3, 2021
7219bff
Use SubsetSubject, return after error, no documentation, add system h…
jtmott-intel Jun 4, 2021
3373830
I guess not documented means not on the supported list
jtmott-intel Jun 4, 2021
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
8 changes: 8 additions & 0 deletions clang/include/clang/Basic/Attr.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1197,6 +1197,14 @@ def SYCLDevice : InheritableAttr {
let Documentation = [SYCLDeviceDocs];
}

def SYCLGlobalVar : InheritableAttr {
let Spellings = [GNU<"sycl_global_var">];
let Subjects = SubjectList<[GlobalVar]>;
let LangOpts = [SYCLIsDevice];
let Documentation = [SYCLGlobalVarDocs];
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Now that this is only usable in system headers, this can go back to being undocumented with a comment about it only being used internally by the SYCL implementation (and the docs can be removed from AttrDocs.td).

let SimpleHandler = 1;
}

def SYCLKernel : InheritableAttr {
let Spellings = [Clang<"sycl_kernel">];
let Subjects = SubjectList<[FunctionTmpl]>;
Expand Down
24 changes: 24 additions & 0 deletions clang/include/clang/Basic/AttrDocs.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3155,6 +3155,30 @@ implicitly inherit this attribute.
}];
}

def SYCLGlobalVarDocs : Documentation {
let Category = DocCatFunction;
let Content = [{
Normally, a SYCL kernel cannot access a global variable, but there are cases
where it is desirable to use a global variable allocated and accessed on a SYCL device. This
attribute is only available to a SYCL device compiler (that is, when passing
``-fsycl-is-device``) and only applies to global variables. It affects semantic
checks to allow use of a marked global within a SYCL kernel.

.. code-block:: c++

#ifdef __SYCL_DEVICE_ONLY__
__attribute__((sycl_global_var)) int Var;
#endif

void F1(cl::sycl::handler& CGH) {
CGH.parallel_for_impl<class TU1>([=] () {
Var = 42; // device code
});
}

}];
}

def RISCVInterruptDocs : Documentation {
let Category = DocCatFunction;
let Heading = "interrupt (RISCV)";
Expand Down
5 changes: 3 additions & 2 deletions clang/lib/Sema/SemaExpr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -223,9 +223,10 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
if (IsRuntimeEvaluated && !IsConst && VD->getStorageClass() == SC_Static)
SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
<< Sema::KernelNonConstStaticDataVariable;
// Non-const globals are allowed for SYCL explicit SIMD.
// Non-const globals are allowed for SYCL explicit SIMD or with the
// SYCLGlobalVar attribute.
else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
VD->hasGlobalStorage())
VD->hasGlobalStorage() && !VD->hasAttr<SYCLGlobalVarAttr>())
SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict)
<< Sema::KernelGlobalVariable;
// ESIMD globals cannot be used in a SYCL context.
Expand Down
1 change: 1 addition & 0 deletions clang/test/Misc/pragma-attribute-supported-attributes-list.test
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -152,6 +152,7 @@
// CHECK-NEXT: ReturnsTwice (SubjectMatchRule_function)
// CHECK-NEXT: SYCLDevice (SubjectMatchRule_function)
// CHECK-NEXT: SYCLDeviceIndirectlyCallable (SubjectMatchRule_function)
// CHECK-NEXT: SYCLGlobalVar (SubjectMatchRule_variable_is_global)
// CHECK-NEXT: SYCLIntelFPGADisableLoopPipelining (SubjectMatchRule_function)
// CHECK-NEXT: SYCLIntelFPGAInitiationInterval (SubjectMatchRule_function)
// CHECK-NEXT: SYCLIntelFPGAMaxConcurrency (SubjectMatchRule_function)
Expand Down
54 changes: 54 additions & 0 deletions clang/test/SemaSYCL/attr-syclglobalvar.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,54 @@
// RUN: %clang_cc1 -fsycl-is-device -verify -fsyntax-only %s

#include "Inputs/sycl.hpp"

__attribute__((sycl_global_var)) int GlobalWithAttribute;

__attribute__((sycl_global_var)) extern int ExternGlobalWithAttribute;

namespace NS {
__attribute__((sycl_global_var)) int NSGlobalWithAttribute;
}

union U {
int InstanceMember;
};

__attribute__((sycl_global_var)) U GlobalUnion;

struct S {
__attribute__((sycl_global_var)) static int StaticMember;

// expected-warning@+1 {{attribute only applies to global variables}}
__attribute__((sycl_global_var)) int InstanceMember;
};
int S::StaticMember = 0;

__attribute__((sycl_global_var)) S GlobalStruct;

__attribute__((sycl_global_var)) static S StaticGlobal;

// expected-error@+1 {{attribute takes no arguments}}
__attribute__((sycl_global_var(42))) int GlobalWithAttributeArg;

int GlobalNoAttribute;

// expected-warning@+1 {{attribute only applies to global variables}}
__attribute__((sycl_global_var)) void F() {
__attribute__((sycl_global_var)) static int StaticLocalVar;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would have expected a diagnostic here as a local variable is not a global variable.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would have expected a diagnostic here as a local variable is not a global variable.

This is an interesting question. SYCL spec applies the "global variable" restriction (that it must be const) to variables of static storage duration.
Here: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:language.restrictions.kernels
and here: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:commonAddressSpace

Given this, I wonder if the attribute could be applied to the file-scope static variables to override the restriction.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I updated the tests and added the anonymous static union.

The anonymous union is not considered a global (doesn't accept the attribute) and isn't allowed in a kernel.

The static local var, surprisingly, is considered a global and accepts the attribute, though still not allowed in a kernel (because I didn't use the attribute's presence to allow that specific case).

These are interesting results. In my changes, I relied on the pre-defined "GlobalVar" subject list. "GlobalVar" is defined earlier as a Var where the code check S->hasGlobalStorage() is true.

I'm stepping through code so I can answer why some of these results "have global storage" and some don't.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Ultimately, though, this isn't a feature we're expecting external users to use. This is meant to be used by us to enable the assert feature. Perhaps there's a different subject list I can use that applies to fewer things?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more. 

  /// Returns true for all variables that do not have local storage.
  ///
  /// This includes all global variables as well as static variables declared
  /// within a function.
  bool hasGlobalStorage() const { return !hasLocalStorage(); }

That's why you get the behavior you're getting (and the name GlobalVar in Attr.td is awful, IMHO).

Ultimately, though, this isn't a feature we're expecting external users to use. This is meant to be used by us to enable the assert feature. Perhaps there's a different subject list I can use that applies to fewer things?

It doesn't much matter what we expect -- once we expose the attribute, users will use the attribute. This is basically giving users a blessed way to opt-in to using global variables whenever they'd like (I'm not convinced this is a good idea, btw) because we found a use case where we want to do it ourselves. We can give this attribute whatever semantics we'd like, so there may be a more restrictive predicate we can use, but without a design document for the feature as such, I don't really know what the right answer is as a reviewer (I can't go off the SYCL 2020 spec because this attribute is intended specifically to do things the SYCL spec doesn't allow).

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Sounds good.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Latest commit expands the test to distinguish between user code and system-include code. Currently, the attribute is still allowed to appear in user code, but it won't have any effect. The attribute will only have an effect if it's used from a system header. To disallow the attribute from even appearing in user code, I'll need to figure out how to do a source location check from within a tablegen subject list.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

On the second thought, here is another example:

// sycl.hpp, system header
template <typename T>
struct Struct {
  __attribute__((sycl_global_var)) static T Variable;
};

// main.cpp, user's application
struct UserTy {
  int X;
  int Y;
};

Struct<UserTy> Global;

Is this use-case allowed?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I would imagine that use case should be allowed -- the declaration of the variable is appropriately marked, the fact that the type comes from a user instantiation shouldn't matter to the attribute's semantics.

Copy link
Contributor Author

@jtmott-intel jtmott-intel Jun 2, 2021 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Latest commit updated with non-simple handler. Also included templates in tests. Struct<UserTy>::StaticMember would work, but Struct<UserTy> Global; would not because you can't use a global in the kernel, and you can't apply the attribute to a global in the user's cpp file (only system headers).


// expected-warning@+1 {{attribute only applies to global variables}}
__attribute__((sycl_global_var)) int Local;

cl::sycl::kernel_single_task<class kernel_name>([=] () {
(void)GlobalWithAttribute;
(void)ExternGlobalWithAttribute;
(void)NS::NSGlobalWithAttribute;
(void)GlobalUnion.InstanceMember;
(void)S::StaticMember;
(void)GlobalStruct.InstanceMember;
(void)StaticGlobal.InstanceMember; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
(void)StaticLocalVar; // expected-error {{SYCL kernel cannot use a non-const static data variable}}
(void)GlobalNoAttribute; // expected-error {{SYCL kernel cannot use a non-const global variable}} expected-note@Inputs/sycl.hpp:* {{called by}}
});
}