Skip to content

[clang] Accept non const static work_group_static variables in templates #20944

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

Draft
Fznamznon wants to merge 1 commit into
base: sycl
Choose a base branch
from
+22 −1
Draft

[clang] Accept non const static work_group_static variables in templates #20944

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
1 change: 1 addition & 0 deletions clang/lib/Sema/SemaDecl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -8197,6 +8197,7 @@ NamedDecl *Sema::ActOnVariableDeclarator(
// constexpr unless their types are decorated with global_variable_allowed
// attribute.
if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) &&
!NewVD->getType()->isDependentType() &&
!SYCL().isTypeDecoratedWithDeclAttribute<SYCLGlobalVariableAllowedAttr>(
NewVD->getType()) &&
!SYCL().isTypeDecoratedWithDeclAttribute<SYCLScopeAttr>(
Expand Down
4 changes: 3 additions & 1 deletion clang/lib/Sema/SemaExpr.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -228,7 +228,8 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
ObjCInterfaceDecl *ClassReceiver,
bool SkipTrailingRequiresClause) {
if (getLangOpts().SYCLIsDevice) {
if (auto VD = dyn_cast<VarDecl>(D)) {
if (auto VD = dyn_cast<VarDecl>(D);
VD && !VD->getType()->isDependentType()) {
Comment on lines -231 to +232
Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

Nice! I forgot that if-with-separate-initializer-and-condition was in C++17 and that we can use it now! Love it!

bool IsConst = VD->getType().isConstant(Context);
bool IsRuntimeEvaluated =
ExprEvalContexts.empty() ||
Expand All @@ -239,6 +240,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> Locs,
if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
VD->getStorageClass() == SC_Static &&
!VD->hasAttr<SYCLGlobalVarAttr>() &&
!VD->getType()->isDependentType() &&
Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

This check is redundant because of the check in the above condition, yes?

Copy link
Contributor Author

@Fznamznon Fznamznon Dec 19, 2025

Choose a reason for hiding this comment

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

Right, I forgot to remove it.

!SemaSYCL::isTypeDecoratedWithDeclAttribute<
SYCLGlobalVariableAllowedAttr>(VD->getType()) &&
!SemaSYCL::isTypeDecoratedWithDeclAttribute<SYCLScopeAttr>(
Expand Down
18 changes: 18 additions & 0 deletions clang/test/SemaSYCL/sycl_wg_scope.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -58,6 +58,8 @@ template <typename T> class [[__sycl_detail__::wg_scope]] B12 { // #B12
B12() = default;
~B12() = default;
T obj;
operator T&() {return obj;}
T *operator&() noexcept { return &obj; }
Comment on lines +61 to +62
Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

It isn't clear to me what these operators are intended to exercise. The conversion operator doesn't appear to be used anywhere.

Copy link
Contributor Author

@Fznamznon Fznamznon Dec 19, 2025

Choose a reason for hiding this comment

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

The ptr method of DependentWrap or ex. MyClass uses the conversion. I can probably make it return G19<T> * instead. Also why 19? there 19 G variables but only 3 classes, lets make it G4

Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

I think ptr() is relying on the T* operator&() declaration. As far as I can tell, the operator T&() declaration isn't exercised. I'm not clear what purpose either of these declarations serve though, couldn't ptr() return &slm.obj directly? Or does the indirection through these operators exercise something more subtle that I'm not picking up on?

And oops, I noticed a g18 declaration and didn't even consider the distinction of variables and classes! Yes, G4 is much better! Or maybe B16 if it will be used to exercise "bad" cases.

};

B12<Valid> b12;
Expand All @@ -84,11 +86,27 @@ struct Wrap {
static G1 g18;
};

template<int Index, typename T>
struct MyClass {
static B12<T> slm;
T *ptr() {
return &slm;
}
};
Comment on lines +89 to +95
Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

Index isn't used. MyClass isn't a very descriptive name; I suggest DependentWrap for consistency above. The B* names appear (mostly) intended to exercise "bad" cases that result in a diagnostic with G* used for "good" cases. If this test is only intended to exercise "good" cases, then I would recommend adding G19 as shown below. However, I think it would be useful to test some "bad" cases as is done with B7 and B12 above (via the Valid and Invalid* cases). In that case, continuing to use B12 seems fine and we should validate that a diagnostic is still issued for, e.g., DependentWrap<InvalidCtor>.

Suggested change
template<int Index, typename T>
struct MyClass {
static B12<T> slm;
T *ptr() {
return &slm;
}
};
template <typename T> class G19 {
T field;
};
template<typename T>
struct DependentWrap {
static G19<T> slm;
T *ptr() {
return &slm;
}
};


template <typename T>
void foo() {
static T m;
}

__attribute__((sycl_device)) void ref_func() {
G1 g19;
static G1 g20;

(void)g16;
(void)g17;
(void)Wrap::g18;
MyClass<1, int> c1;
(void)c1.ptr();
foo<B12<int>>();
Comment on lines +109 to +111
Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

I suggest testing both valid and invalid cases as is done for B7 and B12 elsewhere.

Suggested change
MyClass<1, int> c1;
(void)c1.ptr();
foo<B12<int>>();
DependentWrap<Valid> gdw;
(void)gdw.ptr();
DependentWrap<InvalidCtor> bdw;
(void)bdw.ptr();
foo<B12<Valid>>();
foo<B12<InvalidCtor>>();

This is a tangent, but it could also be useful to test cases of inhibited and deleted default constructors by adding to the set of earlier Invalid* class definitions. I'm not sure if these should trigger the "SYCL work group scope only applies to class with a trivial default constructor" error or perhaps result in some other error (in which case, adding these cases might not be helpful).

struct InhibitedDefaultCtor {
  InhibitedDefaultCtor(const InhibitedDefaultCtor&) = default;
};
struct DeletedDefaultCtor {
  DeletedDefaultCtor() = delete;
};

Copy link
Contributor Author

@Fznamznon Fznamznon Dec 19, 2025 •
edited
Loading

Choose a reason for hiding this comment

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

I think Valid/Invalid classes are intended to test that the attribute cannot be applied to a class which doesn't really make sense for this particular patch, I think?.
Especially combining this with the previous comment #20944 (comment) where G19 doesn't have the attribute makes everything invalid automatically.
What we want to test for this patch is probably that the diagnostic is not lost for a known invalid case, right? Not how the attribute itself works

Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

What I was aiming for (and completely messed up) was to validate that diagnostics are correctly issued when the (non-const) static data member is dependent but is resolved to a type with and without the wg_scope attribute during template instantiation. So, correct, not trying to exercise how the attribute itself works, but rather to ensure that the diagnostic is still issued for a (non-const) static data member originally written in a dependent context that has a type without the wg_scope attribute (in which case, should we also test that the diagnostic is suppressed if the static data member is declared const? I guess that is likely tested elsewhere).

And oops, I missed adding the attribute to G19.

}
Copy link
Contributor

@gmlueck gmlueck Dec 19, 2025

Choose a reason for hiding this comment

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

Will a case like this still be diagnosed as an error?

#include <sycl/sycl.hpp>

template<typename T>
struct MyClass {
  static sycl::marray<T, 2> var;
  int val(int i) {
    return var[i];
  }
};
template<typename T>
sycl::marray<T, 2> MyClass<T>::var;

int main() {
  sycl::queue q;
  sycl::nd_range ndr{{1024}, {16}};
  q.parallel_for(ndr, [=](sycl::nd_item<1> it) {
    MyClass<int> c1;
    c1.val(0);
  }).wait();
}

Here, MyClass declares a static variable which has a dependent type. When MyClass is instantiated with a type that is not decorated with [[__sycl_detail__::wg_scope]] (as in this test), the compiler does diagnose an error. I'm wondering if this will still be diagnosed as an error because the PR comment says "This patch fixes the problem by not diagnosing variables with dependent types".

Copy link
Contributor Author

@Fznamznon Fznamznon Dec 19, 2025

Choose a reason for hiding this comment

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

Yep, it will be diagnosed because once the template is instantiated, the of the variable type is not dependent anymore. The code is analyzed many times by clang before and after instantiation of templates. The problem was coming from the code that is not instantiated yet because the check for the attribute won't succeed.

Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

I'll defer to @Fznamznon for an authoritative reply, but I think the intent is to suppress the diagnostic from being issued in a dependent context. The diagnostic should still be issued for non-dependent contexts (e.g., the instantiated specializations of MyClass). Some of the test suggestions I added are intended to ensure this is the case for classes declared with [[__sycl_detail__::wg_scope]]. I would hope we have sufficient tests to catch regressions similar to what you describe; if not, we should add such tests.

Another tangent: The diagnostic currently issued fails to reference the location in device code that necessitated the code where the error is detected. It would be nice if we could also issue a note that includes that location. See https://godbolt.org/z/68qarnfEW for example; a note referring to line 17 and the instantiation of ct<int>::ptr() would be helpful.

Copy link
Contributor

@tahonermann tahonermann Dec 19, 2025

Choose a reason for hiding this comment

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

Oops, that last example should be valid after this PR. A better example is https://godbolt.org/z/eeacbMosY. Such a note is issued for this example (presumably because the diagnostic is now issued from a non-dependent context), though a note referencing the instantiation of ct2<int>::ptr() is still missing (and there could be a number of such instantiated contexts).

Loading