Skip to content
Merged
Show file tree
Hide file tree
Changes from 3 commits
Commits
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
4 changes: 1 addition & 3 deletions clang/include/clang/Sema/Sema.h
Original file line number Diff line number Diff line change
Expand Up @@ -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
//
Expand Down
3 changes: 1 addition & 2 deletions clang/lib/Sema/SemaExpr.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -249,8 +249,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef<SourceLocation> 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<FunctionDecl>(D)) {
Expand Down
4 changes: 2 additions & 2 deletions clang/lib/Sema/SemaSYCL.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -4132,15 +4132,15 @@ 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");

if (VD->isInvalidDecl() || !VD->hasInit() || !VD->hasGlobalStorage())
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)
Expand Down
9 changes: 8 additions & 1 deletion clang/test/SemaSYCL/Inputs/sycl.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -319,7 +319,14 @@ namespace ext {
namespace oneapi {
namespace experimental {
template <typename T, typename ID = T>
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
Expand Down
26 changes: 26 additions & 0 deletions clang/test/SemaSYCL/spec-const-value-dependent-crash.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,26 @@
// 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<int, class MyInt32Const> SC(default_val);

myQueue.submit([&](sycl::handler &h) {
h.single_task<class kernel_sc>(
[=] {
cl::sycl::ext::oneapi::experimental::spec_constant<int, class MyInt32Const> 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<int, class MyInt32Const>':'sycl::ext::oneapi::experimental::spec_constant<int, MyInt32Const>' 'void ()'