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

Fznamznon · 2025-12-19T14:37:10Z

In general, SYCL doesn't allow non const static variables, but for some extensions like work_group_static it is allowed. This is done by checking presence of an attribute on the type of the variable. However detection of such cases worked incorrectly in template contexts since a dependent type may not yet have the attribute. This patch fixes the problem by not diagnosing variables with dependent types.

tahonermann

The general changes look good. It looks like one of the added checks is redundant. I added some suggestions for the modified test.

tahonermann · 2025-12-19T15:37:53Z

clang/lib/Sema/SemaExpr.cpp

-    if (auto VD = dyn_cast<VarDecl>(D)) {
+    if (auto VD = dyn_cast<VarDecl>(D);
+        VD && !VD->getType()->isDependentType()) {


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

tahonermann · 2025-12-19T15:39:19Z

clang/lib/Sema/SemaExpr.cpp

      if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst &&
          VD->getStorageClass() == SC_Static &&
          !VD->hasAttr<SYCLGlobalVarAttr>() &&
+          !VD->getType()->isDependentType() &&


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

Right, I forgot to remove it.

tahonermann · 2025-12-19T15:52:14Z

clang/test/SemaSYCL/sycl_wg_scope.cpp

+template<int Index, typename T>
+struct MyClass {
+  static B12<T> slm;
+  T *ptr() {
+    return &slm;
+  }
+};


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;

}

};

tahonermann · 2025-12-19T15:57:18Z

clang/test/SemaSYCL/sycl_wg_scope.cpp

+  operator T&() {return obj;}
+  T *operator&() noexcept { return &obj; }


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

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

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.

tahonermann · 2025-12-19T16:11:26Z

clang/test/SemaSYCL/sycl_wg_scope.cpp

+  MyClass<1, int> c1;
+  (void)c1.ptr();
+  foo<B12<int>>();


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; };

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

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.

gmlueck · 2025-12-19T16:15:31Z

clang/test/SemaSYCL/sycl_wg_scope.cpp

+  MyClass<1, int> c1;
+  (void)c1.ptr();
+  foo<B12<int>>();
 }


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".

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.

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.

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).

Fznamznon · 2025-12-19T17:53:44Z

Ah, apparently there is a bug somewhere and these two messages https://godbolt.org/z/Yj7WTzx5j are converted to none by this patch. I kind of expected we would have a test in check-clang for a case that easy but we don't. Let's debug it after the holidays.

tahonermann · 2025-12-19T21:30:54Z

Ah, apparently there is a bug somewhere and these two messages https://godbolt.org/z/Yj7WTzx5j are converted to none by this patch. I kind of expected we would have a test in check-clang for a case that easy but we don't. Let's debug it after the holidays.

Debugging after the holidays sounds great! Enjoy the break (or, if you don't see this until next year, I hope you enjoyed the break)!

There are a couple of interesting things about those particular diagnostics.

The second occurrence doesn't repeat the source code from the file. Switching the order of the calls to foo() doesn't change this, so the missing source reference isn't specific to a call. Perhaps this is just how Clang works when multiple errors referencing the same source location are issued? I don't recall seeing this before.
No notes are issued referencing the location of the calls to foo(); those would be helpful.
Adding const to the declaration of foo()::m suffices to avoid the diagnostics. However, passing, e.g., const Valid as the template argument does not suffice. It seems the diagnostic code must be checking for const as written rather than the type of the static data member.

Fznamznon requested a review from a team as a code owner December 19, 2025 14:37

Fznamznon temporarily deployed to WindowsCILock December 19, 2025 14:38 — with GitHub Actions Inactive

Fznamznon temporarily deployed to WindowsCILock December 19, 2025 15:11 — with GitHub Actions Inactive

tahonermann reviewed Dec 19, 2025

View reviewed changes

gmlueck reviewed Dec 19, 2025

View reviewed changes

Fznamznon marked this pull request as draft December 19, 2025 17:53

		operator T&() {return obj;}
		T *operator&() noexcept { return &obj; }

-  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>>();

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

Are you sure you want to change the base?

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

Uh oh!

Conversation

Fznamznon commented Dec 19, 2025

Uh oh!

tahonermann left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Fznamznon Dec 19, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Fznamznon commented Dec 19, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

tahonermann commented Dec 19, 2025

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Fznamznon Dec 19, 2025 •

edited

Loading

Fznamznon commented Dec 19, 2025 •

edited

Loading