diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 8e71c19ebc9d6..df22cc6f1faa7 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10593,9 +10593,7 @@ class Sema final { Expr *E); bool checkNSReturnsRetainedReturnType(SourceLocation loc, QualType type); - bool checkAllowedSYCLInitializer(VarDecl *VD, - bool CheckValueDependent = false); - + bool checkAllowedSYCLInitializer(VarDecl *VD); //===--------------------------------------------------------------------===// // C++ Coroutines TS // diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 00840d32f1ca0..2b5335904b969 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -249,8 +249,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, // Disallow const statics and globals that are not zero-initialized // or constant-initialized. else if (IsRuntimeEvaluated && IsConst && VD->hasGlobalStorage() && - !VD->isConstexpr() && - !checkAllowedSYCLInitializer(VD, /*CheckValueDependent =*/true)) + !VD->isConstexpr() && !checkAllowedSYCLInitializer(VD)) SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) << Sema::KernelConstStaticVariable; } else if (auto *FDecl = dyn_cast(D)) { diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index fe2f405df22f6..eff50e3a3c649 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -4132,7 +4132,7 @@ void Sema::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, } } -bool Sema::checkAllowedSYCLInitializer(VarDecl *VD, bool CheckValueDependent) { +bool Sema::checkAllowedSYCLInitializer(VarDecl *VD) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); @@ -4140,7 +4140,7 @@ bool Sema::checkAllowedSYCLInitializer(VarDecl *VD, bool CheckValueDependent) { return true; const Expr *Init = VD->getInit(); - bool ValueDependent = CheckValueDependent && Init->isValueDependent(); + bool ValueDependent = Init->isValueDependent(); bool isConstantInit = Init && !ValueDependent && Init->isConstantInitializer(Context, false); if (!VD->isConstexpr() && Init && !ValueDependent && !isConstantInit) diff --git a/clang/test/SemaSYCL/Inputs/sycl.hpp b/clang/test/SemaSYCL/Inputs/sycl.hpp index 90c5f3042cb66..9cb22bce7177b 100644 --- a/clang/test/SemaSYCL/Inputs/sycl.hpp +++ b/clang/test/SemaSYCL/Inputs/sycl.hpp @@ -319,7 +319,14 @@ namespace ext { namespace oneapi { namespace experimental { template -class spec_constant {}; +class spec_constant { +public: + spec_constant() {} + explicit constexpr spec_constant(T defaultVal) : DefaultValue(defaultVal) {} + +private: + T DefaultValue; +}; } // namespace experimental } // namespace oneapi } // namespace ext diff --git a/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp b/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp new file mode 100644 index 0000000000000..7ddb5671add1f --- /dev/null +++ b/clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp @@ -0,0 +1,24 @@ +// RUN: %clang_cc1 -fsycl-is-device -internal-isystem %S/Inputs -ast-dump -sycl-std=2020 %s | FileCheck %s +// This test checks that Clang doesn't crash if a specialization constant is +// value dependent. + +#include "sycl.hpp" +sycl::queue myQueue; + +int main() { + constexpr int default_val = 20; + cl::sycl::ext::oneapi::experimental::spec_constant SC(default_val); + + myQueue.submit([&](sycl::handler &h) { + h.single_task( + [=] { + cl::sycl::ext::oneapi::experimental::spec_constant res = SC; + }); + }); + return 0; +} + +// CHECK: FunctionDecl {{.*}}kernel_sc{{.*}} 'void ()' +// CHECK: VarDecl {{.*}}'(lambda at {{.*}}' +// CHECK-NEXT: InitListExpr {{.*}}'(lambda at {{.*}}' +// CHECK-NEXT: CXXConstructExpr {{.*}}'cl::sycl::ext::oneapi::experimental::spec_constant':'sycl::ext::oneapi::experimental::spec_constant' 'void ()'