[clang] [clang][SYCL] Diagnose reference kernel parameters (PR #192957)
Mariya Podchishchaeva via cfe-commits
cfe-commits at lists.llvm.org
Tue Jun 30 06:15:10 PDT 2026
https://github.com/Fznamznon updated https://github.com/llvm/llvm-project/pull/192957
>From cd2a8069daacb37f781a36c5bd21029de6427cc8 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 20 Apr 2026 04:59:25 -0700
Subject: [PATCH 01/39] [clang][SYCL] Diagnose reference kernel arguments
Per SYCL 2020 spec: Reference types are not trivially copyable, so they may not
be passed as kernel parameters.
This PR adds infrastructure for kernel object visiting and implements
diagnostics for reference kernel parameters.
The infrastructure will be also used for other kernel parameter restrictions and
functional code transformations that will be done in separate PRs.
---
clang/lib/Sema/SemaSYCL.cpp | 278 ++++++++++++++++++
.../SemaSYCL/sycl-kernel-arg-restrictions.cpp | 85 ++++++
2 files changed, 363 insertions(+)
create mode 100644 clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 112a6e4416df2..0d22676aced29 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -294,6 +294,268 @@ void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) {
}
}
+namespace {
+/// A special visitor to visit subobjects within a type, i.e. fields of a
+/// class or elements of an array. Useful for SYCl because in SYCL kernels are
+/// defined via lambda expressions or named callable objects and kernel
+/// parameters are fields of these. These visitors will be used for diagnosing
+/// invalid kernel arugments as well as for functional transformations.
+class SubobjectVisitor {
+ ASTContext &Ctx;
+
+ // These enable handler execution only when previous Handlers succeed.
+ template <typename... Tn>
+ bool handleField(FieldDecl *FD, QualType FDTy, Tn &&...tn) {
+ bool result = true;
+ (void)std::initializer_list<int>{(result = result && tn(FD, FDTy), 0)...};
+ return result;
+ }
+ template <typename... Tn>
+ bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&...tn) {
+ bool result = true;
+ std::initializer_list<int>{(result = result && tn(BD, BDTy), 0)...};
+ return result;
+ }
+
+#define KF_FOR_EACH(FUNC, Item, Qt) \
+ handleField(Item, Qt, ([&](FieldDecl *FD, QualType FDTy) { \
+ return Handlers.FUNC(FD, FDTy); \
+ })...)
+
+ // Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter
+ // the Wrapper structure that we're currently visiting. Owner is the parent
+ // type (which doesn't exist in cases where it is a FieldDecl in the
+ // 'root'), and Wrapper is the current struct being unwrapped.
+ template <typename ParentTy, typename... HandlerTys>
+ void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+ const CXXRecordDecl *Wrapper, QualType RecordTy,
+ HandlerTys &...Handlers) {
+ (void)std::initializer_list<int>{
+ (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...};
+ visitRecordHelper(Wrapper, Wrapper->bases(), Handlers...);
+ visitRecordHelper(Wrapper, Wrapper->fields(), Handlers...);
+ (void)std::initializer_list<int>{
+ (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...};
+ }
+
+ template <typename... HandlerTys>
+ void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+ QualType ArrayTy, HandlerTys &...Handlers) {
+ // TODO add support for simple array visiting, i.e. without entering array
+ // elements.
+ visitComplexArray(Owner, Field, ArrayTy, Handlers...);
+ }
+
+ template <typename ParentTy, typename... HandlerTys>
+ void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
+ const CXXRecordDecl *Wrapper, QualType RecordTy,
+ HandlerTys &...Handlers) {
+ // TODO add support for simple record visiting, i.e. without entering record
+ // fields.
+ visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...);
+ }
+
+ template <typename... HandlerTys>
+ void visitRecordHelper(const CXXRecordDecl *Owner,
+ clang::CXXRecordDecl::base_class_const_range Range,
+ HandlerTys &...Handlers) {
+ for (const auto &Base : Range) {
+ QualType BaseTy = Base.getType();
+ visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy,
+ Handlers...);
+ }
+ }
+
+ template <typename... HandlerTys>
+ void visitRecordHelper(const CXXRecordDecl *Owner, RecordDecl::field_range,
+ HandlerTys &...Handlers) {
+ visitRecordFields(Owner, Handlers...);
+ }
+
+ template <typename... HandlerTys>
+ void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
+ QualType ElementTy, uint64_t Index,
+ HandlerTys &...Handlers) {
+ visitField(Owner, ArrayField, ElementTy, Handlers...);
+ }
+
+ template <typename... HandlerTys>
+ void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
+ QualType ElementTy, uint64_t Index,
+ HandlerTys &...Handlers) {
+ visitArrayElementImpl(Owner, ArrayField, ElementTy, Index, Handlers...);
+ }
+
+ template <typename... HandlerTys>
+ void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field,
+ QualType ArrayTy, HandlerTys &...Handlers) {
+ // Array workflow is:
+ // handleArrayType
+ // enterArray
+ // visitField (same as before, note that The FieldDecl is the of array
+ // itself, not the element)
+ // ... repeat per element, opt-out for duplicates.
+ // leaveArray
+
+ if (!KF_FOR_EACH(handleArrayType, Field, ArrayTy))
+ return;
+
+ const ConstantArrayType *CAT = Ctx.getAsConstantArrayType(ArrayTy);
+ assert(CAT && "Should only be called on constant-size array.");
+ QualType ET = CAT->getElementType();
+ uint64_t ElemCount = CAT->getSize().getZExtValue();
+
+ (void)std::initializer_list<int>{
+ (Handlers.enterArray(Field, ArrayTy, ET), 0)...};
+
+ for (uint64_t Index = 0; Index < ElemCount; ++Index)
+ visitNthArrayElement(Owner, Field, ET, Index, Handlers...);
+
+ (void)std::initializer_list<int>{
+ (Handlers.leaveArray(Field, ArrayTy, ET), 0)...};
+ }
+
+ template <typename... HandlerTys>
+ void visitField(const CXXRecordDecl *Owner, FieldDecl *Field,
+ QualType FieldTy, HandlerTys &...Handlers) {
+ if (FieldTy->isStructureOrClassType()) {
+ if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
+ CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
+ visitRecord(Owner, Field, RD, FieldTy, Handlers...);
+ }
+ } else if (FieldTy->isUnionType())
+ KF_FOR_EACH(handleUnionType, Field, FieldTy);
+ else if (FieldTy->isReferenceType())
+ KF_FOR_EACH(handleReferenceType, Field, FieldTy);
+ else if (FieldTy->isPointerType())
+ KF_FOR_EACH(handlePointerType, Field, FieldTy);
+ else if (FieldTy->isArrayType())
+ visitArray(Owner, Field, FieldTy, Handlers...);
+ else if (FieldTy->isScalarType() || FieldTy->isVectorType())
+ KF_FOR_EACH(handleScalarType, Field, FieldTy);
+ else
+ KF_FOR_EACH(handleOtherType, Field, FieldTy);
+ }
+
+public:
+ SubobjectVisitor(ASTContext &C) : Ctx(C) {}
+
+ template <typename... HandlerTys>
+ void visitRecordBases(const CXXRecordDecl *KernelFunctor,
+ HandlerTys &...Handlers) {
+ visitRecordHelper(KernelFunctor, KernelFunctor->bases(), Handlers...);
+ }
+
+ template <typename... HandlerTys>
+ void visitRecordFields(const CXXRecordDecl *Owner, HandlerTys &...Handlers) {
+ for (const auto Field : Owner->fields())
+ visitField(Owner, Field, Field->getType(), Handlers...);
+ }
+
+#undef KF_FOR_EACH
+};
+
+class SyclKernelFieldHandlerBase {
+public:
+ virtual bool handleStructType(FieldDecl *, QualType) { return true; }
+ virtual bool handleUnionType(FieldDecl *, QualType) { return true; }
+ virtual bool handleReferenceType(FieldDecl *, QualType) { return true; }
+ virtual bool handlePointerType(FieldDecl *, QualType) { return true; }
+ virtual bool handleArrayType(FieldDecl *, QualType) { return true; }
+ virtual bool handleScalarType(FieldDecl *, QualType) { return true; }
+ // Most handlers shouldn't be handling this, just the field checker.
+ virtual bool handleOtherType(FieldDecl *, QualType) { return true; }
+
+ virtual bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
+ return true;
+ }
+ virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
+ return true;
+ }
+ virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
+ QualType) {
+ return true;
+ }
+ virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
+ QualType) {
+ return true;
+ }
+ // The following are used for stepping through array elements.
+ virtual bool enterArray(FieldDecl *, QualType, QualType) { return true; }
+ virtual bool leaveArray(FieldDecl *, QualType, QualType) { return true; }
+
+ virtual ~SyclKernelFieldHandlerBase() = default;
+};
+
+// A class to act as the direct base for all the SYCL Kernel related
+// tasks that contains a reference to Sema (and potentially any other
+// universally required data).
+class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase {
+protected:
+ SemaSYCL &SemaSYCLRef;
+ SyclKernelFieldHandler(SemaSYCL &S) : SemaSYCLRef(S) {}
+};
+
+// A type to check the validity of all of the argument types.
+class SyclKernelFieldChecker : public SyclKernelFieldHandler {
+ bool IsInvalid = false;
+
+ bool checkArrayType(QualType Ty, SourceLocation Loc) {
+ if (Ty->isArrayType()) {
+ if (const auto *CAT =
+ SemaSYCLRef.getASTContext().getAsConstantArrayType(Ty)) {
+ QualType ET = CAT->getElementType();
+ return checkArrayType(ET, Loc);
+ }
+ return SemaSYCLRef.Diag(Loc, diag::err_bad_kernel_param_type) << Ty;
+ }
+ return false;
+ }
+
+public:
+ /// Constructor for the SyclKernelFieldChecker
+ /// \param S The SemaSYCL reference used for diagnostics and context.
+ explicit SyclKernelFieldChecker(SemaSYCL &S) : SyclKernelFieldHandler(S) {}
+ bool isValid() { return !IsInvalid; }
+
+ bool handleReferenceType(FieldDecl *FD, QualType FieldTy) final {
+ SemaSYCLRef.Diag(FD->getLocation(), diag::err_bad_kernel_param_type)
+ << FieldTy;
+ IsInvalid = true;
+ return isValid();
+ }
+
+ bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
+ while (FieldTy->isAnyPointerType()) {
+ FieldTy = QualType{FieldTy->getPointeeOrArrayElementType(), 0};
+ if (FieldTy->isVariableArrayType()) {
+ SemaSYCLRef.Diag(FD->getLocation(), diag::err_bad_kernel_param_type)
+ << FieldTy;
+ IsInvalid = true;
+ break;
+ }
+ }
+ return isValid();
+ }
+
+ bool handleOtherType(FieldDecl *FD, QualType FieldTy) final {
+ SemaSYCLRef.Diag(FD->getLocation(), diag::err_bad_kernel_param_type)
+ << FieldTy;
+ IsInvalid = true;
+ return isValid();
+ }
+
+ bool handleStructType(FieldDecl *, QualType FieldTy) final {
+ return isValid();
+ }
+
+ bool handleArrayType(FieldDecl *FD, QualType FieldTy) final {
+ IsInvalid |= checkArrayType(FieldTy, FD->getLocation());
+ return isValid();
+ }
+};
+} // namespace
+
void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
// Ensure that all attributes present on the declaration are consistent
// and warn about any redundant ones.
@@ -665,6 +927,20 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
+bool verifyKernelArguments(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
+ SyclKernelFieldChecker FieldChecker(SemaSYCLRef);
+ SubobjectVisitor Visitor{SemaSYCLRef.getASTContext()};
+ for (auto Param : FD->parameters()) {
+ if (Param->getType()->isRecordType()) {
+ const CXXRecordDecl *ObjRecord = Param->getType()->getAsCXXRecordDecl();
+ assert(ObjRecord && "object is expected");
+ Visitor.visitRecordBases(ObjRecord, FieldChecker);
+ Visitor.visitRecordFields(ObjRecord, FieldChecker);
+ }
+ }
+ return !FieldChecker.isValid();
+}
+
} // unnamed namespace
StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
@@ -690,6 +966,8 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
getASTContext().getSYCLKernelInfo(SKEPAttr->getKernelName());
assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
"SYCL kernel name conflict");
+ if (verifyKernelArguments(FD, *this))
+ return StmtError();
// Build the outline of the synthesized device entry point function.
OutlinedFunctionDecl *OFD =
diff --git a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
new file mode 100644
index 0000000000000..6025f89ff12c4
--- /dev/null
+++ b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
@@ -0,0 +1,85 @@
+// RUN: %clang_cc1 -triple x86_64-linux-gnu -std=c++17 -fsyntax-only -Wno-vla-cxx-extension -fsycl-is-host -verify %s
+// RUN: %clang_cc1 -triple spirv64 -std=c++17 -fsyntax-only -Wno-vla-cxx-extension -fsycl-is-device -verify %s
+
+// A unique kernel name type is required for each declared kernel entry point.
+template<int, int = 0> struct KN;
+
+// A generic kernel launch function.
+template<typename KNT, typename... Ts>
+void sycl_kernel_launch(const char *, Ts...) {}
+
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T) {}
+
+struct Empty {};
+
+struct S {
+ int a;
+ int &b; //expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
+};
+
+void fooarr(int (&arr)[5]) {
+}
+
+template <typename T> class Callable {
+ T data; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+public:
+ Callable(T d) : data(d) {}
+ void operator()() {
+ }
+};
+
+void refCases(int AS) {
+ int p = 0;
+ double q = 0;
+ float s = 0;
+ kernel_single_task<class KN<1>>( // expected-note {{requested here}}
+ [
+ // expected-error at +1 {{'int &' cannot be used as the type of a kernel parameter}}
+ &p, q,
+ // expected-error at +1 {{'float &' cannot be used as the type of a kernel parameter}}
+ &s] {
+ (void)q;
+ (void)p;
+ (void)s;
+ });
+
+ auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+ S Str {p, p};
+ kernel_single_task<class KN<2>>( // expected-note {{requested here}}
+ [=] {
+ (void)L;
+ (void)Str; // no error because fail for L already
+ });
+
+ kernel_single_task<class KN<3>>( // expected-note {{requested here}}
+ [=] {
+ (void)Str;
+ });
+
+ S arr[2] = {Str, Str};
+ kernel_single_task<class KN<4>>( // expected-note {{requested here}}
+ [=] {
+ (void)arr;
+ });
+ int arr1[AS];
+ kernel_single_task<class KN<5>>( // expected-note {{requested here}}
+ [&] {
+ (void)arr1; // expected-error {{'int (&)[AS]' cannot be used as the type of a kernel parameter}}
+ });
+ auto a = &arr1;
+ kernel_single_task<class KN<6>>( // expected-note {{requested here}}
+ [=] {
+ (void)a; // expected-error {{'int[AS]' cannot be used as the type of a kernel parameter}}
+ });
+ int arrayints[5] = {0};
+ kernel_single_task<class KN<7>>( // expected-note {{requested here}}
+ [&] {
+ fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
+ });
+ kernel_single_task<class KN<8>>(Callable<int&>{p}); // expected-note {{requested here}}
+ kernel_single_task<class KN<9>>(Callable<int>{p});
+}
+
>From 6a0361a7069999d7032b83ed2748ca8636ac59e9 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 8 May 2026 08:52:01 -0700
Subject: [PATCH 02/39] Redo the visitor as CRTP, put it outside of SemaSYCL
---
clang/include/clang/AST/SubobjectVisitor.h | 140 ++++++++
clang/lib/Sema/SemaSYCL.cpp | 307 +++---------------
.../SemaSYCL/sycl-kernel-arg-restrictions.cpp | 33 +-
3 files changed, 196 insertions(+), 284 deletions(-)
create mode 100644 clang/include/clang/AST/SubobjectVisitor.h
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
new file mode 100644
index 0000000000000..ee9b70d6710db
--- /dev/null
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -0,0 +1,140 @@
+//===---------- SubobjectVisitor.h - Subobject Visitor ----------*- C++ -*-===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+//
+// This file defines the SubobjectVisitor interface, which recursively
+// traverses subobjects within a type.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef LLVM_CLANG_AST_SUBOBJECTVISITOR_H
+#define LLVM_CLANG_AST_SUBOBJECTVISITOR_H
+
+#include "clang/AST/Type.h"
+
+namespace clang {
+template <typename Derived> class SubobjectVisitor {
+ ASTContext &Ctx;
+
+ public:
+ SubobjectVisitor(ASTContext &Ctx) : Ctx(Ctx) {}
+ /// Return a reference to the derived class.
+ Derived &getDerived() { return *static_cast<Derived *>(this); }
+
+ bool enterRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
+ return getDerived().visitRecord(Record, Parent);
+ }
+
+ bool leaveRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool enterUnion(CXXRecordDecl *Record, FieldDecl *Parent) {
+ return getDerived().visitUnion(Record, Parent);
+ }
+
+ bool enterArray(QualType ArrayTy, FieldDecl *Parent) {
+ return getDerived().visitArrayTy(ArrayTy, Parent);
+ }
+
+ bool visitRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool visitUnion(CXXRecordDecl *Record, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool visitReferenceType(QualType Ty, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool visitPointerType(QualType Ty, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool visitScalarType(QualType Ty, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool visitOtherType(QualType Ty, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool visitArrayTy(QualType Ty, FieldDecl *Parent) {
+ return true;
+ }
+
+ bool traverseRecord(CXXRecordDecl *Record) {
+ FieldDecl* Parent = nullptr;
+ getDerived().traverseRecord(Record, Parent);
+ return true;
+ }
+
+ bool traverseRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
+ getDerived().enterRecord(Record, Parent);
+
+ for (const auto &Base : Record->bases()) {
+ QualType BaseTy = Base.getType();
+ getDerived().traverseType(BaseTy, Parent);
+ }
+ for (const auto Field : Record->fields()) {
+ QualType FieldTy = Field->getType();
+ getDerived().traverseType(FieldTy, Field);
+ }
+
+ getDerived().leaveRecord(Record, Parent);
+ return true;
+ }
+
+ bool traverseUnion(CXXRecordDecl *Record, FieldDecl *Parent) {
+ getDerived().enterUnion(Record, Parent);
+
+ for (const auto Field : Record->fields()) {
+ QualType FieldTy = Field->getType();
+ getDerived().traverseType(FieldTy, Field);
+ }
+ return true;
+ }
+
+ bool traverseArray(QualType Ty, FieldDecl *Parent) {
+ getDerived().enterArray(Ty, Parent);
+ const ConstantArrayType *CAT = Ctx.getAsConstantArrayType(Ty);
+ if (!CAT)
+ return true;
+
+ QualType ET = CAT->getElementType();
+ uint64_t ElemCount = CAT->getSize().getZExtValue();
+ for (uint64_t Index = 0; Index < ElemCount; ++Index)
+ getDerived().traverseType(ET, Parent);
+
+ return true;
+ }
+
+ bool traverseType(QualType Ty, FieldDecl *Parent) {
+ if (Ty->isStructureOrClassType()) {
+ CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
+ getDerived().traverseRecord(RD, Parent);
+ } else if (Ty->isUnionType()) {
+ CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
+ getDerived().traverseUnion(RD, Parent);
+ } else if (Ty->isReferenceType())
+ getDerived().visitReferenceType(Ty, Parent);
+ else if (Ty->isPointerType())
+ getDerived().visitPointerType(Ty, Parent);
+ else if (Ty->isArrayType())
+ getDerived().traverseArray(Ty, Parent);
+ else if (Ty->isScalarType() || Ty->isVectorType())
+ getDerived().visitScalarType(Ty, Parent);
+ else
+ getDerived().visitOtherType(Ty, Parent);
+ return true;
+ }
+
+};
+} // end namespace clang
+#endif // LLVM_CLANG_AST_SUBOBJECTVISITOR
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 0d22676aced29..ce16f25fb34f9 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -11,6 +11,8 @@
#include "clang/Sema/SemaSYCL.h"
#include "TreeTransform.h"
#include "clang/AST/Mangle.h"
+#include "clang/AST/RecursiveASTVisitor.h"
+#include "clang/AST/SubobjectVisitor.h"
#include "clang/AST/SYCLKernelInfo.h"
#include "clang/AST/StmtSYCL.h"
#include "clang/AST/TypeOrdering.h"
@@ -294,268 +296,6 @@ void SemaSYCL::CheckSYCLExternalFunctionDecl(FunctionDecl *FD) {
}
}
-namespace {
-/// A special visitor to visit subobjects within a type, i.e. fields of a
-/// class or elements of an array. Useful for SYCl because in SYCL kernels are
-/// defined via lambda expressions or named callable objects and kernel
-/// parameters are fields of these. These visitors will be used for diagnosing
-/// invalid kernel arugments as well as for functional transformations.
-class SubobjectVisitor {
- ASTContext &Ctx;
-
- // These enable handler execution only when previous Handlers succeed.
- template <typename... Tn>
- bool handleField(FieldDecl *FD, QualType FDTy, Tn &&...tn) {
- bool result = true;
- (void)std::initializer_list<int>{(result = result && tn(FD, FDTy), 0)...};
- return result;
- }
- template <typename... Tn>
- bool handleField(const CXXBaseSpecifier &BD, QualType BDTy, Tn &&...tn) {
- bool result = true;
- std::initializer_list<int>{(result = result && tn(BD, BDTy), 0)...};
- return result;
- }
-
-#define KF_FOR_EACH(FUNC, Item, Qt) \
- handleField(Item, Qt, ([&](FieldDecl *FD, QualType FDTy) { \
- return Handlers.FUNC(FD, FDTy); \
- })...)
-
- // Parent contains the FieldDecl or CXXBaseSpecifier that was used to enter
- // the Wrapper structure that we're currently visiting. Owner is the parent
- // type (which doesn't exist in cases where it is a FieldDecl in the
- // 'root'), and Wrapper is the current struct being unwrapped.
- template <typename ParentTy, typename... HandlerTys>
- void visitComplexRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
- const CXXRecordDecl *Wrapper, QualType RecordTy,
- HandlerTys &...Handlers) {
- (void)std::initializer_list<int>{
- (Handlers.enterStruct(Owner, Parent, RecordTy), 0)...};
- visitRecordHelper(Wrapper, Wrapper->bases(), Handlers...);
- visitRecordHelper(Wrapper, Wrapper->fields(), Handlers...);
- (void)std::initializer_list<int>{
- (Handlers.leaveStruct(Owner, Parent, RecordTy), 0)...};
- }
-
- template <typename... HandlerTys>
- void visitArray(const CXXRecordDecl *Owner, FieldDecl *Field,
- QualType ArrayTy, HandlerTys &...Handlers) {
- // TODO add support for simple array visiting, i.e. without entering array
- // elements.
- visitComplexArray(Owner, Field, ArrayTy, Handlers...);
- }
-
- template <typename ParentTy, typename... HandlerTys>
- void visitRecord(const CXXRecordDecl *Owner, ParentTy &Parent,
- const CXXRecordDecl *Wrapper, QualType RecordTy,
- HandlerTys &...Handlers) {
- // TODO add support for simple record visiting, i.e. without entering record
- // fields.
- visitComplexRecord(Owner, Parent, Wrapper, RecordTy, Handlers...);
- }
-
- template <typename... HandlerTys>
- void visitRecordHelper(const CXXRecordDecl *Owner,
- clang::CXXRecordDecl::base_class_const_range Range,
- HandlerTys &...Handlers) {
- for (const auto &Base : Range) {
- QualType BaseTy = Base.getType();
- visitRecord(Owner, Base, BaseTy->getAsCXXRecordDecl(), BaseTy,
- Handlers...);
- }
- }
-
- template <typename... HandlerTys>
- void visitRecordHelper(const CXXRecordDecl *Owner, RecordDecl::field_range,
- HandlerTys &...Handlers) {
- visitRecordFields(Owner, Handlers...);
- }
-
- template <typename... HandlerTys>
- void visitArrayElementImpl(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
- QualType ElementTy, uint64_t Index,
- HandlerTys &...Handlers) {
- visitField(Owner, ArrayField, ElementTy, Handlers...);
- }
-
- template <typename... HandlerTys>
- void visitNthArrayElement(const CXXRecordDecl *Owner, FieldDecl *ArrayField,
- QualType ElementTy, uint64_t Index,
- HandlerTys &...Handlers) {
- visitArrayElementImpl(Owner, ArrayField, ElementTy, Index, Handlers...);
- }
-
- template <typename... HandlerTys>
- void visitComplexArray(const CXXRecordDecl *Owner, FieldDecl *Field,
- QualType ArrayTy, HandlerTys &...Handlers) {
- // Array workflow is:
- // handleArrayType
- // enterArray
- // visitField (same as before, note that The FieldDecl is the of array
- // itself, not the element)
- // ... repeat per element, opt-out for duplicates.
- // leaveArray
-
- if (!KF_FOR_EACH(handleArrayType, Field, ArrayTy))
- return;
-
- const ConstantArrayType *CAT = Ctx.getAsConstantArrayType(ArrayTy);
- assert(CAT && "Should only be called on constant-size array.");
- QualType ET = CAT->getElementType();
- uint64_t ElemCount = CAT->getSize().getZExtValue();
-
- (void)std::initializer_list<int>{
- (Handlers.enterArray(Field, ArrayTy, ET), 0)...};
-
- for (uint64_t Index = 0; Index < ElemCount; ++Index)
- visitNthArrayElement(Owner, Field, ET, Index, Handlers...);
-
- (void)std::initializer_list<int>{
- (Handlers.leaveArray(Field, ArrayTy, ET), 0)...};
- }
-
- template <typename... HandlerTys>
- void visitField(const CXXRecordDecl *Owner, FieldDecl *Field,
- QualType FieldTy, HandlerTys &...Handlers) {
- if (FieldTy->isStructureOrClassType()) {
- if (KF_FOR_EACH(handleStructType, Field, FieldTy)) {
- CXXRecordDecl *RD = FieldTy->getAsCXXRecordDecl();
- visitRecord(Owner, Field, RD, FieldTy, Handlers...);
- }
- } else if (FieldTy->isUnionType())
- KF_FOR_EACH(handleUnionType, Field, FieldTy);
- else if (FieldTy->isReferenceType())
- KF_FOR_EACH(handleReferenceType, Field, FieldTy);
- else if (FieldTy->isPointerType())
- KF_FOR_EACH(handlePointerType, Field, FieldTy);
- else if (FieldTy->isArrayType())
- visitArray(Owner, Field, FieldTy, Handlers...);
- else if (FieldTy->isScalarType() || FieldTy->isVectorType())
- KF_FOR_EACH(handleScalarType, Field, FieldTy);
- else
- KF_FOR_EACH(handleOtherType, Field, FieldTy);
- }
-
-public:
- SubobjectVisitor(ASTContext &C) : Ctx(C) {}
-
- template <typename... HandlerTys>
- void visitRecordBases(const CXXRecordDecl *KernelFunctor,
- HandlerTys &...Handlers) {
- visitRecordHelper(KernelFunctor, KernelFunctor->bases(), Handlers...);
- }
-
- template <typename... HandlerTys>
- void visitRecordFields(const CXXRecordDecl *Owner, HandlerTys &...Handlers) {
- for (const auto Field : Owner->fields())
- visitField(Owner, Field, Field->getType(), Handlers...);
- }
-
-#undef KF_FOR_EACH
-};
-
-class SyclKernelFieldHandlerBase {
-public:
- virtual bool handleStructType(FieldDecl *, QualType) { return true; }
- virtual bool handleUnionType(FieldDecl *, QualType) { return true; }
- virtual bool handleReferenceType(FieldDecl *, QualType) { return true; }
- virtual bool handlePointerType(FieldDecl *, QualType) { return true; }
- virtual bool handleArrayType(FieldDecl *, QualType) { return true; }
- virtual bool handleScalarType(FieldDecl *, QualType) { return true; }
- // Most handlers shouldn't be handling this, just the field checker.
- virtual bool handleOtherType(FieldDecl *, QualType) { return true; }
-
- virtual bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
- return true;
- }
- virtual bool leaveStruct(const CXXRecordDecl *, FieldDecl *, QualType) {
- return true;
- }
- virtual bool enterStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
- QualType) {
- return true;
- }
- virtual bool leaveStruct(const CXXRecordDecl *, const CXXBaseSpecifier &,
- QualType) {
- return true;
- }
- // The following are used for stepping through array elements.
- virtual bool enterArray(FieldDecl *, QualType, QualType) { return true; }
- virtual bool leaveArray(FieldDecl *, QualType, QualType) { return true; }
-
- virtual ~SyclKernelFieldHandlerBase() = default;
-};
-
-// A class to act as the direct base for all the SYCL Kernel related
-// tasks that contains a reference to Sema (and potentially any other
-// universally required data).
-class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase {
-protected:
- SemaSYCL &SemaSYCLRef;
- SyclKernelFieldHandler(SemaSYCL &S) : SemaSYCLRef(S) {}
-};
-
-// A type to check the validity of all of the argument types.
-class SyclKernelFieldChecker : public SyclKernelFieldHandler {
- bool IsInvalid = false;
-
- bool checkArrayType(QualType Ty, SourceLocation Loc) {
- if (Ty->isArrayType()) {
- if (const auto *CAT =
- SemaSYCLRef.getASTContext().getAsConstantArrayType(Ty)) {
- QualType ET = CAT->getElementType();
- return checkArrayType(ET, Loc);
- }
- return SemaSYCLRef.Diag(Loc, diag::err_bad_kernel_param_type) << Ty;
- }
- return false;
- }
-
-public:
- /// Constructor for the SyclKernelFieldChecker
- /// \param S The SemaSYCL reference used for diagnostics and context.
- explicit SyclKernelFieldChecker(SemaSYCL &S) : SyclKernelFieldHandler(S) {}
- bool isValid() { return !IsInvalid; }
-
- bool handleReferenceType(FieldDecl *FD, QualType FieldTy) final {
- SemaSYCLRef.Diag(FD->getLocation(), diag::err_bad_kernel_param_type)
- << FieldTy;
- IsInvalid = true;
- return isValid();
- }
-
- bool handlePointerType(FieldDecl *FD, QualType FieldTy) final {
- while (FieldTy->isAnyPointerType()) {
- FieldTy = QualType{FieldTy->getPointeeOrArrayElementType(), 0};
- if (FieldTy->isVariableArrayType()) {
- SemaSYCLRef.Diag(FD->getLocation(), diag::err_bad_kernel_param_type)
- << FieldTy;
- IsInvalid = true;
- break;
- }
- }
- return isValid();
- }
-
- bool handleOtherType(FieldDecl *FD, QualType FieldTy) final {
- SemaSYCLRef.Diag(FD->getLocation(), diag::err_bad_kernel_param_type)
- << FieldTy;
- IsInvalid = true;
- return isValid();
- }
-
- bool handleStructType(FieldDecl *, QualType FieldTy) final {
- return isValid();
- }
-
- bool handleArrayType(FieldDecl *FD, QualType FieldTy) final {
- IsInvalid |= checkArrayType(FieldTy, FD->getLocation());
- return isValid();
- }
-};
-} // namespace
-
void SemaSYCL::CheckSYCLEntryPointFunctionDecl(FunctionDecl *FD) {
// Ensure that all attributes present on the declaration are consistent
// and warn about any redundant ones.
@@ -927,18 +667,49 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
+class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
+ SemaSYCL &SemaSYCLRef;
+ bool IsValid = true;
+ SmallVector<const CXXRecordDecl *, 4> History;
+
+public:
+ KernelArgsChecker(SemaSYCL &SR, SourceLocation Loc)
+ : SubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
+ SemaSYCLRef(SR) {}
+ bool enterRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
+ History.push_back(Record);
+ return getDerived().visitRecord(Record, Parent);
+ }
+
+ bool visitReferenceType(QualType Ty, FieldDecl *Parent) {
+ SemaSYCLRef.Diag(Parent->getLocation(), diag::err_bad_kernel_param_type)
+ << Ty;
+ for (auto &ParentRecord : History)
+ SemaSYCLRef.Diag(ParentRecord->getLocation(),
+ diag::note_within_field_of_type)
+ << ParentRecord;
+ IsValid = false;
+ return true;
+ }
+
+ bool leaveRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
+ History.pop_back();
+ return true;
+ }
+
+ bool isInvalid() { return !IsValid; }
+};
+
bool verifyKernelArguments(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
- SyclKernelFieldChecker FieldChecker(SemaSYCLRef);
- SubobjectVisitor Visitor{SemaSYCLRef.getASTContext()};
+ KernelArgsChecker KAC(SemaSYCLRef, FD->getLocation());
for (auto Param : FD->parameters()) {
if (Param->getType()->isRecordType()) {
- const CXXRecordDecl *ObjRecord = Param->getType()->getAsCXXRecordDecl();
+ CXXRecordDecl *ObjRecord = Param->getType()->getAsCXXRecordDecl();
assert(ObjRecord && "object is expected");
- Visitor.visitRecordBases(ObjRecord, FieldChecker);
- Visitor.visitRecordFields(ObjRecord, FieldChecker);
+ KAC.traverseRecord(ObjRecord);
}
}
- return !FieldChecker.isValid();
+ return KAC.isInvalid();
}
} // unnamed namespace
diff --git a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
index 6025f89ff12c4..2df4dadaf293a 100644
--- a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
@@ -13,17 +13,15 @@ template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
void kernel_single_task(T) {}
-struct Empty {};
-
-struct S {
+struct S { // expected-note 4{{within field of type 'S' declared here}}
int a;
- int &b; //expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
+ int &b; //expected-error 4{{'int &' cannot be used as the type of a kernel parameter}}
};
void fooarr(int (&arr)[5]) {
}
-template <typename T> class Callable {
+template <typename T> class Callable { // expected-note {{within field of type 'Callable<int &>' declared here}}
T data; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
public:
Callable(T d) : data(d) {}
@@ -31,12 +29,18 @@ template <typename T> class Callable {
}
};
+class Derived1 : Callable<int> { // expected-note {{within field of type 'Derived1' declared here}}
+ int &a; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+public:
+ Derived1(int d, int &b) : Callable<int>(d), a(b) {}
+};
+
void refCases(int AS) {
int p = 0;
double q = 0;
float s = 0;
kernel_single_task<class KN<1>>( // expected-note {{requested here}}
- [
+ [ // expected-note2{{within field of type}}
// expected-error at +1 {{'int &' cannot be used as the type of a kernel parameter}}
&p, q,
// expected-error at +1 {{'float &' cannot be used as the type of a kernel parameter}}
@@ -47,39 +51,36 @@ void refCases(int AS) {
});
auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+ // expected-note at -1 {{within field of type}}
S Str {p, p};
kernel_single_task<class KN<2>>( // expected-note {{requested here}}
- [=] {
+ [=] { // expected-note 2{{within field of type}}
(void)L;
(void)Str; // no error because fail for L already
});
kernel_single_task<class KN<3>>( // expected-note {{requested here}}
- [=] {
+ [=] { // // expected-note {{within field of type}}
(void)Str;
});
S arr[2] = {Str, Str};
kernel_single_task<class KN<4>>( // expected-note {{requested here}}
- [=] {
+ [=] { // expected-note 2{{within field of type}}
(void)arr;
});
int arr1[AS];
kernel_single_task<class KN<5>>( // expected-note {{requested here}}
- [&] {
+ [&] { // expected-note {{within field of type}}
(void)arr1; // expected-error {{'int (&)[AS]' cannot be used as the type of a kernel parameter}}
});
- auto a = &arr1;
- kernel_single_task<class KN<6>>( // expected-note {{requested here}}
- [=] {
- (void)a; // expected-error {{'int[AS]' cannot be used as the type of a kernel parameter}}
- });
int arrayints[5] = {0};
kernel_single_task<class KN<7>>( // expected-note {{requested here}}
- [&] {
+ [&] { // expected-note {{within field of type}}
fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
});
kernel_single_task<class KN<8>>(Callable<int&>{p}); // expected-note {{requested here}}
kernel_single_task<class KN<9>>(Callable<int>{p});
+ kernel_single_task<class KN<10>>(Derived1{p, p}); // expected-note {{requested here}}
}
>From b6c566d9a91fd94644e12db383ccf7dc6a3eb23f Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Wed, 20 May 2026 10:29:50 -0700
Subject: [PATCH 03/39] WIP, test doesn't pass yet
---
clang/include/clang/AST/SubobjectVisitor.h | 181 ++++++++----------
.../clang/Basic/DiagnosticSemaKinds.td | 2 +
clang/lib/Sema/SemaSYCL.cpp | 50 +++--
3 files changed, 122 insertions(+), 111 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index ee9b70d6710db..62c8e6ec65c02 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -17,124 +17,113 @@
#include "clang/AST/Type.h"
namespace clang {
-template <typename Derived> class SubobjectVisitor {
- ASTContext &Ctx;
- public:
- SubobjectVisitor(ASTContext &Ctx) : Ctx(Ctx) {}
- /// Return a reference to the derived class.
- Derived &getDerived() { return *static_cast<Derived *>(this); }
-
- bool enterRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
- return getDerived().visitRecord(Record, Parent);
- }
+#define DISPATCH(CLASS) \
+ return static_cast<Derived *>(this)->visit##CLASS( \
+ static_cast<const CLASS *>(T))
- bool leaveRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
- return true;
- }
+template <template <typename> class Ptr, typename Derived>
+class SubobjectVisitorBase {
+ ASTContext &Ctx;
+ template <typename Class> using ptr_t = typename Ptr<Class>::type;
+ template <typename Class>
+ using non_ptr_t = typename std::remove_pointer<ptr_t<Class>>::type;
- bool enterUnion(CXXRecordDecl *Record, FieldDecl *Parent) {
- return getDerived().visitUnion(Record, Parent);
- }
+public:
+ SubobjectVisitorBase(ASTContext &Ctx) : Ctx(Ctx) {}
- bool enterArray(QualType ArrayTy, FieldDecl *Parent) {
- return getDerived().visitArrayTy(ArrayTy, Parent);
- }
-
- bool visitRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
- return true;
- }
+ /// Return a reference to the derived class.
+ Derived &getDerived() { return *static_cast<Derived *>(this); }
- bool visitUnion(CXXRecordDecl *Record, FieldDecl *Parent) {
- return true;
- }
+ void visit(QualType QT) {
+ // If the type is an array, visit its element type. Separate traversal of
+ // arrays is not needed because the array will be encountered as a
+ // FieldDecl.
- bool visitReferenceType(QualType Ty, FieldDecl *Parent) {
- return true;
- }
+ if (QT->isArrayType()) {
+ QualType ElTy =
+ cast<ConstantArrayType>(Ctx.getAsArrayType(QT))->getElementType();
+ visit(ElTy);
+ return;
+ }
- bool visitPointerType(QualType Ty, FieldDecl *Parent) {
- return true;
- }
+ if (ptr_t<RecordDecl> RD = QT->getAsRecordDecl()) {
+ traverseRecord(RD);
+ return;
+ }
- bool visitScalarType(QualType Ty, FieldDecl *Parent) {
- return true;
+ visitGenericType(QT.getTypePtr());
}
- bool visitOtherType(QualType Ty, FieldDecl *Parent) {
- return true;
+ void traverseRecord(ptr_t<RecordDecl> RD) {
+ if (ptr_t<CXXRecordDecl> CRD = dyn_cast<CXXRecordDecl>(RD)) {
+ for (non_ptr_t<CXXBaseSpecifier> BS : CRD->bases()) {
+ if (!static_cast<Derived *>(this)->visitBaseSpecifierPre(BS))
+ continue;
+ visit(BS.getType());
+ static_cast<Derived *>(this)->visitBaseSpecifierPost(BS);
+ }
+ }
+ for (ptr_t<FieldDecl> FD : RD->fields()) {
+ if (!static_cast<Derived *>(this)->visitFieldDeclPre(FD))
+ continue;
+ visit(FD->getType());
+ static_cast<Derived *>(this)->visitFieldDeclPost(FD);
+ }
}
- bool visitArrayTy(QualType Ty, FieldDecl *Parent) {
- return true;
- }
+ // Default base class specifier pre-order visitor.
+ bool visitBaseSpecifierPre(non_ptr_t<CXXBaseSpecifier> BS) { return true; }
- bool traverseRecord(CXXRecordDecl *Record) {
- FieldDecl* Parent = nullptr;
- getDerived().traverseRecord(Record, Parent);
- return true;
- }
+ // Default base class specifier post-order visitor.
+ void visitBaseSpecifierPost(non_ptr_t<CXXBaseSpecifier> BS) {}
- bool traverseRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
- getDerived().enterRecord(Record, Parent);
+ // Default field pre-order visitor.
+ bool visitFieldDeclPre(ptr_t<FieldDecl> FD) { return true; }
- for (const auto &Base : Record->bases()) {
- QualType BaseTy = Base.getType();
- getDerived().traverseType(BaseTy, Parent);
+ // Default field post-order visitor.
+ void visitFieldDeclPost(ptr_t<FieldDecl> FD) {}
+ bool visitGenericType(const Type *T) {
+ // Top switch stmt: dispatch to VisitFooType for each FooType.
+ switch (T->getTypeClass()) {
+#define ABSTRACT_TYPE(CLASS, PARENT)
+#define TYPE(CLASS, PARENT) \
+ case Type::CLASS: \
+ DISPATCH(CLASS##Type);
+#include "clang/AST/TypeNodes.inc"
}
- for (const auto Field : Record->fields()) {
- QualType FieldTy = Field->getType();
- getDerived().traverseType(FieldTy, Field);
- }
-
- getDerived().leaveRecord(Record, Parent);
- return true;
+ llvm_unreachable("Unknown type class!");
}
- bool traverseUnion(CXXRecordDecl *Record, FieldDecl *Parent) {
- getDerived().enterUnion(Record, Parent);
+ // If the implementation chooses not to implement a certain visit method, fall
+ // back on superclass.
+#define TYPE(CLASS, PARENT) \
+ bool visit##CLASS##Type(const CLASS##Type *T) { DISPATCH(PARENT); }
+#include "clang/AST/TypeNodes.inc"
- for (const auto Field : Record->fields()) {
- QualType FieldTy = Field->getType();
- getDerived().traverseType(FieldTy, Field);
- }
- return true;
- }
-
- bool traverseArray(QualType Ty, FieldDecl *Parent) {
- getDerived().enterArray(Ty, Parent);
- const ConstantArrayType *CAT = Ctx.getAsConstantArrayType(Ty);
- if (!CAT)
- return true;
+ /// Method called if \c ImpClass doesn't provide specific handler
+ /// for some type class.
+ bool visitType(const Type *) { return true; }
+};
- QualType ET = CAT->getElementType();
- uint64_t ElemCount = CAT->getSize().getZExtValue();
- for (uint64_t Index = 0; Index < ElemCount; ++Index)
- getDerived().traverseType(ET, Parent);
+template <typename Derived>
+class SubobjectVisitor
+ : public SubobjectVisitorBase<std::add_pointer, Derived> {
+public:
+ SubobjectVisitor(ASTContext &Ctx)
+ : SubobjectVisitorBase<std::add_pointer, Derived>(Ctx) {}
+};
- return true;
- }
+template <typename Derived>
+class ConstSubobjectVisitor
+ : public SubobjectVisitorBase<llvm::make_const_ptr, Derived> {
+public:
+ ConstSubobjectVisitor(ASTContext &Ctx)
+ : SubobjectVisitorBase<std::add_pointer, Derived>(Ctx) {}
+};
- bool traverseType(QualType Ty, FieldDecl *Parent) {
- if (Ty->isStructureOrClassType()) {
- CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
- getDerived().traverseRecord(RD, Parent);
- } else if (Ty->isUnionType()) {
- CXXRecordDecl *RD = Ty->getAsCXXRecordDecl();
- getDerived().traverseUnion(RD, Parent);
- } else if (Ty->isReferenceType())
- getDerived().visitReferenceType(Ty, Parent);
- else if (Ty->isPointerType())
- getDerived().visitPointerType(Ty, Parent);
- else if (Ty->isArrayType())
- getDerived().traverseArray(Ty, Parent);
- else if (Ty->isScalarType() || Ty->isVectorType())
- getDerived().visitScalarType(Ty, Parent);
- else
- getDerived().visitOtherType(Ty, Parent);
- return true;
- }
+#undef DISPATCH
-};
} // end namespace clang
+
#endif // LLVM_CLANG_AST_SUBOBJECTVISITOR
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 9cd3a5b2fc049..0f05e798b92b3 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11804,6 +11804,8 @@ def err_record_with_pointers_kernel_param : Error<
"%select{struct|union}0 kernel parameters may not contain pointers">;
def note_within_field_of_type : Note<
"within field of type %0 declared here">;
+def note_within_base_of_type : Note<
+ "within base of type %1 declared here">;
def note_illegal_field_declared_here : Note<
"field of illegal %select{type|pointer type}0 %1 declared here">;
def err_opencl_type_struct_or_union_field : Error<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index ce16f25fb34f9..469bdef168509 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -11,7 +11,6 @@
#include "clang/Sema/SemaSYCL.h"
#include "TreeTransform.h"
#include "clang/AST/Mangle.h"
-#include "clang/AST/RecursiveASTVisitor.h"
#include "clang/AST/SubobjectVisitor.h"
#include "clang/AST/SYCLKernelInfo.h"
#include "clang/AST/StmtSYCL.h"
@@ -670,33 +669,54 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
- SmallVector<const CXXRecordDecl *, 4> History;
+ using SubobjectAccess = llvm::PointerUnion<CXXBaseSpecifier *, FieldDecl *>;
+ SmallVector<SubobjectAccess, 4> SubobjectAccessPath;
public:
KernelArgsChecker(SemaSYCL &SR, SourceLocation Loc)
: SubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
SemaSYCLRef(SR) {}
- bool enterRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
- History.push_back(Record);
- return getDerived().visitRecord(Record, Parent);
+
+ bool visitBaseSpecifierPre(CXXBaseSpecifier BS) {
+ SubobjectAccessPath.push_back(&BS);
+ return true;
}
- bool visitReferenceType(QualType Ty, FieldDecl *Parent) {
- SemaSYCLRef.Diag(Parent->getLocation(), diag::err_bad_kernel_param_type)
- << Ty;
- for (auto &ParentRecord : History)
- SemaSYCLRef.Diag(ParentRecord->getLocation(),
- diag::note_within_field_of_type)
- << ParentRecord;
- IsValid = false;
+ bool visitFieldDeclPre(FieldDecl *FD) {
+ SubobjectAccessPath.push_back(FD);
return true;
}
- bool leaveRecord(CXXRecordDecl *Record, FieldDecl *Parent) {
- History.pop_back();
+ bool visitReferenceType(const ReferenceType *RT) {
+
+ // Reference cannot be a base, so just assume we came via a FieldDecl.
+ auto *DirectParent = SubobjectAccessPath.back().get<FieldDecl *>();
+ SemaSYCLRef.Diag(DirectParent->getLocation(),
+ diag::err_bad_kernel_param_type)
+ << DirectParent->getType();
+
+ for (auto Parent : SubobjectAccessPath) {
+ if (auto *FD = Parent.dyn_cast<FieldDecl *>()) {
+ SemaSYCLRef.Diag(FD->getParent()->getLocation(),
+ diag::note_within_field_of_type)
+ << FD->getParent();
+ } else {
+ auto *BS = Parent.get<CXXBaseSpecifier *>();
+ CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
+ assert(RD);
+ SemaSYCLRef.Diag(RD->getLocation(), diag::note_within_base_of_type)
+ << RD;
+ }
+ }
+ IsValid = false;
return true;
}
+ void visitFieldDeclPost(FieldDecl *FD) { SubobjectAccessPath.pop_back(); }
+ void visitBaseSpecifierPost(CXXBaseSpecifier BS) {
+ SubobjectAccessPath.pop_back();
+ }
+
bool isInvalid() { return !IsValid; }
};
>From b06133dc57b14e534928c0c3598690706fa09c4f Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 21 May 2026 04:27:29 -0700
Subject: [PATCH 04/39] Make the test pass
---
clang/include/clang/AST/SubobjectVisitor.h | 2 +-
clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp | 8 ++++----
2 files changed, 5 insertions(+), 5 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 62c8e6ec65c02..0bf70c2a8b6af 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -52,7 +52,7 @@ class SubobjectVisitorBase {
return;
}
- visitGenericType(QT.getTypePtr());
+ visitGenericType(QT.getCanonicalType().getTypePtr());
}
void traverseRecord(ptr_t<RecordDecl> RD) {
diff --git a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
index 2df4dadaf293a..f727b416953f0 100644
--- a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
@@ -13,9 +13,9 @@ template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
void kernel_single_task(T) {}
-struct S { // expected-note 4{{within field of type 'S' declared here}}
+struct S { // expected-note 3{{within field of type 'S' declared here}}
int a;
- int &b; //expected-error 4{{'int &' cannot be used as the type of a kernel parameter}}
+ int &b; //expected-error 3{{'int &' cannot be used as the type of a kernel parameter}}
};
void fooarr(int (&arr)[5]) {
@@ -52,7 +52,7 @@ void refCases(int AS) {
auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
// expected-note at -1 {{within field of type}}
- S Str {p, p};
+ S Str {p, p};
kernel_single_task<class KN<2>>( // expected-note {{requested here}}
[=] { // expected-note 2{{within field of type}}
(void)L;
@@ -66,7 +66,7 @@ void refCases(int AS) {
S arr[2] = {Str, Str};
kernel_single_task<class KN<4>>( // expected-note {{requested here}}
- [=] { // expected-note 2{{within field of type}}
+ [=] { // expected-note {{within field of type}}
(void)arr;
});
int arr1[AS];
>From 90a7172e4845d63c0216b16f8f95ea5f21694e28 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 21 May 2026 06:03:23 -0700
Subject: [PATCH 05/39] Fix base class history
---
clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +-
clang/lib/Sema/SemaSYCL.cpp | 2 +-
clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp | 11 +++++++++--
3 files changed, 11 insertions(+), 4 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 0f05e798b92b3..7d70827341f93 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11805,7 +11805,7 @@ def err_record_with_pointers_kernel_param : Error<
def note_within_field_of_type : Note<
"within field of type %0 declared here">;
def note_within_base_of_type : Note<
- "within base of type %1 declared here">;
+ "within base of type %0 declared here">;
def note_illegal_field_declared_here : Note<
"field of illegal %select{type|pointer type}0 %1 declared here">;
def err_opencl_type_struct_or_union_field : Error<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 469bdef168509..388a471a7d9e5 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -704,7 +704,7 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
auto *BS = Parent.get<CXXBaseSpecifier *>();
CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
assert(RD);
- SemaSYCLRef.Diag(RD->getLocation(), diag::note_within_base_of_type)
+ SemaSYCLRef.Diag(BS->getBeginLoc(), diag::note_within_base_of_type)
<< RD;
}
}
diff --git a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
index f727b416953f0..47ca74d1642a3 100644
--- a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
@@ -21,8 +21,8 @@ struct S { // expected-note 3{{within field of type 'S' declared here}}
void fooarr(int (&arr)[5]) {
}
-template <typename T> class Callable { // expected-note {{within field of type 'Callable<int &>' declared here}}
- T data; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+template <typename T> class Callable { // expected-note 2{{within field of type 'Callable<int &>' declared here}}
+ T data; // expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
public:
Callable(T d) : data(d) {}
void operator()() {
@@ -35,6 +35,12 @@ class Derived1 : Callable<int> { // expected-note {{within field of type 'Derive
Derived1(int d, int &b) : Callable<int>(d), a(b) {}
};
+class Derived2 : Callable<int&> { // expected-note {{within base of type 'Callable<int &>' declared here}}
+ int a;
+public:
+ Derived2(int d, int &b) : Callable<int&>(b), a(d) {}
+};
+
void refCases(int AS) {
int p = 0;
double q = 0;
@@ -82,5 +88,6 @@ void refCases(int AS) {
kernel_single_task<class KN<8>>(Callable<int&>{p}); // expected-note {{requested here}}
kernel_single_task<class KN<9>>(Callable<int>{p});
kernel_single_task<class KN<10>>(Derived1{p, p}); // expected-note {{requested here}}
+ kernel_single_task<class KN<11>>(Derived2{p, p}); // expected-note {{requested here}}
}
>From afc70df320b3051473b507c678cf8d785cbe2c9e Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 21 May 2026 06:08:44 -0700
Subject: [PATCH 06/39] Fix the warning
---
clang/lib/Sema/SemaSYCL.cpp | 5 ++---
1 file changed, 2 insertions(+), 3 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 388a471a7d9e5..503fb9e33f27b 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -688,9 +688,8 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
}
bool visitReferenceType(const ReferenceType *RT) {
-
// Reference cannot be a base, so just assume we came via a FieldDecl.
- auto *DirectParent = SubobjectAccessPath.back().get<FieldDecl *>();
+ auto *DirectParent = cast<FieldDecl *>(SubobjectAccessPath.back());
SemaSYCLRef.Diag(DirectParent->getLocation(),
diag::err_bad_kernel_param_type)
<< DirectParent->getType();
@@ -701,7 +700,7 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
diag::note_within_field_of_type)
<< FD->getParent();
} else {
- auto *BS = Parent.get<CXXBaseSpecifier *>();
+ auto *BS = cast<CXXBaseSpecifier *>(Parent);
CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
assert(RD);
SemaSYCLRef.Diag(BS->getBeginLoc(), diag::note_within_base_of_type)
>From e60ade531053ae0b3662e8fb6edb741e5ac902a9 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 21 May 2026 06:15:25 -0700
Subject: [PATCH 07/39] Visit, not traverse
---
clang/lib/Sema/SemaSYCL.cpp | 7 ++-----
1 file changed, 2 insertions(+), 5 deletions(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 503fb9e33f27b..6d2b3447868ce 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -722,11 +722,8 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
bool verifyKernelArguments(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
KernelArgsChecker KAC(SemaSYCLRef, FD->getLocation());
for (auto Param : FD->parameters()) {
- if (Param->getType()->isRecordType()) {
- CXXRecordDecl *ObjRecord = Param->getType()->getAsCXXRecordDecl();
- assert(ObjRecord && "object is expected");
- KAC.traverseRecord(ObjRecord);
- }
+ if (Param->getType()->isRecordType())
+ KAC.visit(Param->getType());
}
return KAC.isInvalid();
}
>From aa45cd19fded6e5298b288bee6e217cef2ce6649 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 21 May 2026 06:16:08 -0700
Subject: [PATCH 08/39] Make format happy
---
clang/lib/Sema/SemaSYCL.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 6d2b3447868ce..c4a925baa710f 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -11,9 +11,9 @@
#include "clang/Sema/SemaSYCL.h"
#include "TreeTransform.h"
#include "clang/AST/Mangle.h"
-#include "clang/AST/SubobjectVisitor.h"
#include "clang/AST/SYCLKernelInfo.h"
#include "clang/AST/StmtSYCL.h"
+#include "clang/AST/SubobjectVisitor.h"
#include "clang/AST/TypeOrdering.h"
#include "clang/Basic/Diagnostic.h"
#include "clang/Sema/Attr.h"
>From 74f2460cc0da5379ddec79d1624f0305da600803 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 22 May 2026 07:06:29 -0700
Subject: [PATCH 09/39] Use getDerived(), eliminate non_ptr_t
---
clang/include/clang/AST/SubobjectVisitor.h | 27 ++++++++++------------
clang/lib/Sema/SemaSYCL.cpp | 6 ++---
2 files changed, 15 insertions(+), 18 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 0bf70c2a8b6af..22ad184b70b52 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -26,8 +26,6 @@ template <template <typename> class Ptr, typename Derived>
class SubobjectVisitorBase {
ASTContext &Ctx;
template <typename Class> using ptr_t = typename Ptr<Class>::type;
- template <typename Class>
- using non_ptr_t = typename std::remove_pointer<ptr_t<Class>>::type;
public:
SubobjectVisitorBase(ASTContext &Ctx) : Ctx(Ctx) {}
@@ -43,40 +41,39 @@ class SubobjectVisitorBase {
if (QT->isArrayType()) {
QualType ElTy =
cast<ConstantArrayType>(Ctx.getAsArrayType(QT))->getElementType();
- visit(ElTy);
+ getDerived().visit(ElTy);
return;
}
if (ptr_t<RecordDecl> RD = QT->getAsRecordDecl()) {
- traverseRecord(RD);
+ getDerived().traverseRecord(RD);
return;
}
-
- visitGenericType(QT.getCanonicalType().getTypePtr());
+ getDerived().visitGenericType(QT.getCanonicalType().getTypePtr());
}
void traverseRecord(ptr_t<RecordDecl> RD) {
if (ptr_t<CXXRecordDecl> CRD = dyn_cast<CXXRecordDecl>(RD)) {
- for (non_ptr_t<CXXBaseSpecifier> BS : CRD->bases()) {
- if (!static_cast<Derived *>(this)->visitBaseSpecifierPre(BS))
+ for (CXXBaseSpecifier& BS : CRD->bases()) {
+ if (!getDerived().visitBaseSpecifierPre(&BS))
continue;
- visit(BS.getType());
- static_cast<Derived *>(this)->visitBaseSpecifierPost(BS);
+ getDerived().visit(BS.getType());
+ getDerived().visitBaseSpecifierPost(&BS);
}
}
for (ptr_t<FieldDecl> FD : RD->fields()) {
- if (!static_cast<Derived *>(this)->visitFieldDeclPre(FD))
+ if (!getDerived().visitFieldDeclPre(FD))
continue;
- visit(FD->getType());
- static_cast<Derived *>(this)->visitFieldDeclPost(FD);
+ getDerived().visit(FD->getType());
+ getDerived().visitFieldDeclPost(FD);
}
}
// Default base class specifier pre-order visitor.
- bool visitBaseSpecifierPre(non_ptr_t<CXXBaseSpecifier> BS) { return true; }
+ bool visitBaseSpecifierPre(ptr_t<CXXBaseSpecifier> BS) { return true; }
// Default base class specifier post-order visitor.
- void visitBaseSpecifierPost(non_ptr_t<CXXBaseSpecifier> BS) {}
+ void visitBaseSpecifierPost(ptr_t<CXXBaseSpecifier> BS) {}
// Default field pre-order visitor.
bool visitFieldDeclPre(ptr_t<FieldDecl> FD) { return true; }
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c4a925baa710f..7c56bc82e3c10 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -677,8 +677,8 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
: SubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
SemaSYCLRef(SR) {}
- bool visitBaseSpecifierPre(CXXBaseSpecifier BS) {
- SubobjectAccessPath.push_back(&BS);
+ bool visitBaseSpecifierPre(CXXBaseSpecifier *BS) {
+ SubobjectAccessPath.push_back(BS);
return true;
}
@@ -712,7 +712,7 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
}
void visitFieldDeclPost(FieldDecl *FD) { SubobjectAccessPath.pop_back(); }
- void visitBaseSpecifierPost(CXXBaseSpecifier BS) {
+ void visitBaseSpecifierPost(CXXBaseSpecifier *BS) {
SubobjectAccessPath.pop_back();
}
>From 4622ad836a7f88859b3c6e5bfca0d5fc8892ca5e Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 22 May 2026 09:20:49 -0700
Subject: [PATCH 10/39] Remove "concrete type" visitation
---
clang/include/clang/AST/SubobjectVisitor.h | 28 -------
clang/lib/Sema/SemaSYCL.cpp | 95 +++++++++++++++-------
2 files changed, 64 insertions(+), 59 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 22ad184b70b52..116ecc6692eb9 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -18,10 +18,6 @@
namespace clang {
-#define DISPATCH(CLASS) \
- return static_cast<Derived *>(this)->visit##CLASS( \
- static_cast<const CLASS *>(T))
-
template <template <typename> class Ptr, typename Derived>
class SubobjectVisitorBase {
ASTContext &Ctx;
@@ -49,7 +45,6 @@ class SubobjectVisitorBase {
getDerived().traverseRecord(RD);
return;
}
- getDerived().visitGenericType(QT.getCanonicalType().getTypePtr());
}
void traverseRecord(ptr_t<RecordDecl> RD) {
@@ -80,27 +75,6 @@ class SubobjectVisitorBase {
// Default field post-order visitor.
void visitFieldDeclPost(ptr_t<FieldDecl> FD) {}
- bool visitGenericType(const Type *T) {
- // Top switch stmt: dispatch to VisitFooType for each FooType.
- switch (T->getTypeClass()) {
-#define ABSTRACT_TYPE(CLASS, PARENT)
-#define TYPE(CLASS, PARENT) \
- case Type::CLASS: \
- DISPATCH(CLASS##Type);
-#include "clang/AST/TypeNodes.inc"
- }
- llvm_unreachable("Unknown type class!");
- }
-
- // If the implementation chooses not to implement a certain visit method, fall
- // back on superclass.
-#define TYPE(CLASS, PARENT) \
- bool visit##CLASS##Type(const CLASS##Type *T) { DISPATCH(PARENT); }
-#include "clang/AST/TypeNodes.inc"
-
- /// Method called if \c ImpClass doesn't provide specific handler
- /// for some type class.
- bool visitType(const Type *) { return true; }
};
template <typename Derived>
@@ -119,8 +93,6 @@ class ConstSubobjectVisitor
: SubobjectVisitorBase<std::add_pointer, Derived>(Ctx) {}
};
-#undef DISPATCH
-
} // end namespace clang
#endif // LLVM_CLANG_AST_SUBOBJECTVISITOR
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 7c56bc82e3c10..c9eadc26d41bf 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -669,51 +669,86 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
- using SubobjectAccess = llvm::PointerUnion<CXXBaseSpecifier *, FieldDecl *>;
- SmallVector<SubobjectAccess, 4> SubobjectAccessPath;
+ using ObjectAccess =
+ llvm::PointerUnion<ParmVarDecl *, CXXBaseSpecifier *, FieldDecl *>;
+ SmallVector<ObjectAccess, 4> ObjectAccessPath;
+
+ void emitObjectAccessPathNotes() {
+ for (auto Parent : ObjectAccessPath) {
+ if (auto *FD = Parent.dyn_cast<FieldDecl *>()) {
+ SemaSYCLRef.Diag(FD->getParent()->getLocation(),
+ diag::note_within_field_of_type)
+ << FD->getParent();
+ } else if (auto *BS = Parent.dyn_cast<CXXBaseSpecifier *>()) {
+ CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
+ assert(RD);
+ SemaSYCLRef.Diag(BS->getBeginLoc(), diag::note_within_base_of_type)
+ << RD;
+ } else {
+ // Nothing to emit for ParmVarDecl since its location just points to
+ // skep-attributed function template.
+ assert(isa<ParmVarDecl *>(Parent));
+ }
+ }
+ }
public:
KernelArgsChecker(SemaSYCL &SR, SourceLocation Loc)
: SubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
SemaSYCLRef(SR) {}
+ void checkParameter(ParmVarDecl *PVD) {
+ ObjectAccessPath.push_back(PVD);
+ // Check the immediate type of the parameter.
+ if (checkType(PVD->getType())) {
+ // If type checking wasn't short circuited, visit subobjects to check
+ // them.
+ visit(PVD->getType());
+ }
+ ObjectAccessPath.pop_back();
+ // We either visited everything or we stopped visiting due to invalid kernel
+ // argument type.
+ assert(ObjectAccessPath.empty());
+ }
+
bool visitBaseSpecifierPre(CXXBaseSpecifier *BS) {
- SubobjectAccessPath.push_back(BS);
- return true;
+ ObjectAccessPath.push_back(BS);
+ return checkType(BS->getType());
}
bool visitFieldDeclPre(FieldDecl *FD) {
- SubobjectAccessPath.push_back(FD);
- return true;
+ ObjectAccessPath.push_back(FD);
+ return checkType(FD->getType());
}
- bool visitReferenceType(const ReferenceType *RT) {
- // Reference cannot be a base, so just assume we came via a FieldDecl.
- auto *DirectParent = cast<FieldDecl *>(SubobjectAccessPath.back());
- SemaSYCLRef.Diag(DirectParent->getLocation(),
- diag::err_bad_kernel_param_type)
- << DirectParent->getType();
-
- for (auto Parent : SubobjectAccessPath) {
- if (auto *FD = Parent.dyn_cast<FieldDecl *>()) {
- SemaSYCLRef.Diag(FD->getParent()->getLocation(),
- diag::note_within_field_of_type)
- << FD->getParent();
- } else {
- auto *BS = cast<CXXBaseSpecifier *>(Parent);
- CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
- assert(RD);
- SemaSYCLRef.Diag(BS->getBeginLoc(), diag::note_within_base_of_type)
- << RD;
+ // Returns true if subobjects should be visited and false otherwise.
+ bool checkType(QualType Ty) {
+ if (Ty->isReferenceType()) {
+ auto DirectParent = ObjectAccessPath.back();
+ // Reference cannot be a base, so just assume we came via a FieldDecl.
+ if (isa<ParmVarDecl *>(DirectParent)) {
+ // If reference is a kernel argument, there is nothing to do. We allow
+ // references in direct kernel arguments for better performance of the
+ // host code and we eliminate them when building actual kernel.
+ return true;
}
+
+ auto *DirectFieldParent = cast<FieldDecl *>(DirectParent);
+ SemaSYCLRef.Diag(DirectFieldParent->getLocation(),
+ diag::err_bad_kernel_param_type)
+ << DirectFieldParent->getType();
+ emitObjectAccessPathNotes();
+
+ IsValid = false;
+ ObjectAccessPath.pop_back();
+ return false;
}
- IsValid = false;
return true;
}
- void visitFieldDeclPost(FieldDecl *FD) { SubobjectAccessPath.pop_back(); }
+ void visitFieldDeclPost(FieldDecl *FD) { ObjectAccessPath.pop_back(); }
void visitBaseSpecifierPost(CXXBaseSpecifier *BS) {
- SubobjectAccessPath.pop_back();
+ ObjectAccessPath.pop_back();
}
bool isInvalid() { return !IsValid; }
@@ -721,10 +756,8 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
bool verifyKernelArguments(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
KernelArgsChecker KAC(SemaSYCLRef, FD->getLocation());
- for (auto Param : FD->parameters()) {
- if (Param->getType()->isRecordType())
- KAC.visit(Param->getType());
- }
+ for (auto Param : FD->parameters())
+ KAC.checkParameter(Param);
return KAC.isInvalid();
}
>From 46db74989e52a7646bf9714cfa9782b3b5609265 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 26 May 2026 02:43:22 -0700
Subject: [PATCH 11/39] Fix constvisitor
---
clang/include/clang/AST/SubobjectVisitor.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 116ecc6692eb9..d56d905454a5d 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -90,7 +90,7 @@ class ConstSubobjectVisitor
: public SubobjectVisitorBase<llvm::make_const_ptr, Derived> {
public:
ConstSubobjectVisitor(ASTContext &Ctx)
- : SubobjectVisitorBase<std::add_pointer, Derived>(Ctx) {}
+ : SubobjectVisitorBase<llvm::make_const_ptr, Derived>(Ctx) {}
};
} // end namespace clang
>From dc7b5a9f41c62f767f49b28ff1226b649104a220 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 26 May 2026 02:54:56 -0700
Subject: [PATCH 12/39] Enter post visitation even if pre visitation failed
---
clang/include/clang/AST/SubobjectVisitor.h | 10 ++++------
clang/lib/Sema/SemaSYCL.cpp | 3 ---
2 files changed, 4 insertions(+), 9 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index d56d905454a5d..bb36b8e217aec 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -50,16 +50,14 @@ class SubobjectVisitorBase {
void traverseRecord(ptr_t<RecordDecl> RD) {
if (ptr_t<CXXRecordDecl> CRD = dyn_cast<CXXRecordDecl>(RD)) {
for (CXXBaseSpecifier& BS : CRD->bases()) {
- if (!getDerived().visitBaseSpecifierPre(&BS))
- continue;
- getDerived().visit(BS.getType());
+ if (getDerived().visitBaseSpecifierPre(&BS))
+ getDerived().visit(BS.getType());
getDerived().visitBaseSpecifierPost(&BS);
}
}
for (ptr_t<FieldDecl> FD : RD->fields()) {
- if (!getDerived().visitFieldDeclPre(FD))
- continue;
- getDerived().visit(FD->getType());
+ if (getDerived().visitFieldDeclPre(FD))
+ getDerived().visit(FD->getType());
getDerived().visitFieldDeclPost(FD);
}
}
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index c9eadc26d41bf..58ab6537a462d 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -706,8 +706,6 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
visit(PVD->getType());
}
ObjectAccessPath.pop_back();
- // We either visited everything or we stopped visiting due to invalid kernel
- // argument type.
assert(ObjectAccessPath.empty());
}
@@ -740,7 +738,6 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
emitObjectAccessPathNotes();
IsValid = false;
- ObjectAccessPath.pop_back();
return false;
}
return true;
>From 627848eb6f7ed5267f881db4ac07cf2d8f1a7f33 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 26 May 2026 03:06:26 -0700
Subject: [PATCH 13/39] Const it
---
clang/include/clang/AST/SubobjectVisitor.h | 4 +++-
clang/lib/Sema/SemaSYCL.cpp | 27 +++++++++++-----------
2 files changed, 17 insertions(+), 14 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index bb36b8e217aec..ffd14b95d7dec 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -22,6 +22,8 @@ template <template <typename> class Ptr, typename Derived>
class SubobjectVisitorBase {
ASTContext &Ctx;
template <typename Class> using ptr_t = typename Ptr<Class>::type;
+ template <typename Class>
+ using non_ptr_t = typename std::remove_pointer<ptr_t<Class>>::type;
public:
SubobjectVisitorBase(ASTContext &Ctx) : Ctx(Ctx) {}
@@ -49,7 +51,7 @@ class SubobjectVisitorBase {
void traverseRecord(ptr_t<RecordDecl> RD) {
if (ptr_t<CXXRecordDecl> CRD = dyn_cast<CXXRecordDecl>(RD)) {
- for (CXXBaseSpecifier& BS : CRD->bases()) {
+ for (non_ptr_t<CXXBaseSpecifier&> BS : CRD->bases()) {
if (getDerived().visitBaseSpecifierPre(&BS))
getDerived().visit(BS.getType());
getDerived().visitBaseSpecifierPost(&BS);
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 58ab6537a462d..a39c29cd45798 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -666,20 +666,21 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
-class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
+class KernelArgsChecker : public ConstSubobjectVisitor<KernelArgsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
using ObjectAccess =
- llvm::PointerUnion<ParmVarDecl *, CXXBaseSpecifier *, FieldDecl *>;
+ llvm::PointerUnion<const ParmVarDecl *, const CXXBaseSpecifier *,
+ const FieldDecl *>;
SmallVector<ObjectAccess, 4> ObjectAccessPath;
void emitObjectAccessPathNotes() {
for (auto Parent : ObjectAccessPath) {
- if (auto *FD = Parent.dyn_cast<FieldDecl *>()) {
+ if (auto *FD = Parent.dyn_cast<const FieldDecl *>()) {
SemaSYCLRef.Diag(FD->getParent()->getLocation(),
diag::note_within_field_of_type)
<< FD->getParent();
- } else if (auto *BS = Parent.dyn_cast<CXXBaseSpecifier *>()) {
+ } else if (auto *BS = Parent.dyn_cast<const CXXBaseSpecifier *>()) {
CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
assert(RD);
SemaSYCLRef.Diag(BS->getBeginLoc(), diag::note_within_base_of_type)
@@ -687,17 +688,17 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
} else {
// Nothing to emit for ParmVarDecl since its location just points to
// skep-attributed function template.
- assert(isa<ParmVarDecl *>(Parent));
+ assert(isa<const ParmVarDecl *>(Parent));
}
}
}
public:
KernelArgsChecker(SemaSYCL &SR, SourceLocation Loc)
- : SubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
+ : ConstSubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
SemaSYCLRef(SR) {}
- void checkParameter(ParmVarDecl *PVD) {
+ void checkParameter(const ParmVarDecl *PVD) {
ObjectAccessPath.push_back(PVD);
// Check the immediate type of the parameter.
if (checkType(PVD->getType())) {
@@ -709,12 +710,12 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
assert(ObjectAccessPath.empty());
}
- bool visitBaseSpecifierPre(CXXBaseSpecifier *BS) {
+ bool visitBaseSpecifierPre(const CXXBaseSpecifier *BS) {
ObjectAccessPath.push_back(BS);
return checkType(BS->getType());
}
- bool visitFieldDeclPre(FieldDecl *FD) {
+ bool visitFieldDeclPre(const FieldDecl *FD) {
ObjectAccessPath.push_back(FD);
return checkType(FD->getType());
}
@@ -724,14 +725,14 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
if (Ty->isReferenceType()) {
auto DirectParent = ObjectAccessPath.back();
// Reference cannot be a base, so just assume we came via a FieldDecl.
- if (isa<ParmVarDecl *>(DirectParent)) {
+ if (isa<const ParmVarDecl *>(DirectParent)) {
// If reference is a kernel argument, there is nothing to do. We allow
// references in direct kernel arguments for better performance of the
// host code and we eliminate them when building actual kernel.
return true;
}
- auto *DirectFieldParent = cast<FieldDecl *>(DirectParent);
+ auto *DirectFieldParent = cast<const FieldDecl *>(DirectParent);
SemaSYCLRef.Diag(DirectFieldParent->getLocation(),
diag::err_bad_kernel_param_type)
<< DirectFieldParent->getType();
@@ -743,8 +744,8 @@ class KernelArgsChecker : public SubobjectVisitor<KernelArgsChecker> {
return true;
}
- void visitFieldDeclPost(FieldDecl *FD) { ObjectAccessPath.pop_back(); }
- void visitBaseSpecifierPost(CXXBaseSpecifier *BS) {
+ void visitFieldDeclPost(const FieldDecl *FD) { ObjectAccessPath.pop_back(); }
+ void visitBaseSpecifierPost(const CXXBaseSpecifier *BS) {
ObjectAccessPath.pop_back();
}
>From 3730ffce79e368b43c39ee037634c2dd30d9d7da Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 26 May 2026 06:16:08 -0700
Subject: [PATCH 14/39] Organize the test
---
.../SemaSYCL/sycl-kernel-arg-restrictions.cpp | 89 ++++++++++++-------
1 file changed, 59 insertions(+), 30 deletions(-)
diff --git a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
index 47ca74d1642a3..2f7681de9c74e 100644
--- a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
@@ -13,35 +13,9 @@ template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
void kernel_single_task(T) {}
-struct S { // expected-note 3{{within field of type 'S' declared here}}
- int a;
- int &b; //expected-error 3{{'int &' cannot be used as the type of a kernel parameter}}
-};
-
-void fooarr(int (&arr)[5]) {
-}
-
-template <typename T> class Callable { // expected-note 2{{within field of type 'Callable<int &>' declared here}}
- T data; // expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
-public:
- Callable(T d) : data(d) {}
- void operator()() {
- }
-};
-
-class Derived1 : Callable<int> { // expected-note {{within field of type 'Derived1' declared here}}
- int &a; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
-public:
- Derived1(int d, int &b) : Callable<int>(d), a(b) {}
-};
-
-class Derived2 : Callable<int&> { // expected-note {{within base of type 'Callable<int &>' declared here}}
- int a;
-public:
- Derived2(int d, int &b) : Callable<int&>(b), a(d) {}
-};
-
-void refCases(int AS) {
+// Check that reference captures of kernel that defined as lambda are diagnosed.
+namespace badref1 {
+void test() {
int p = 0;
double q = 0;
float s = 0;
@@ -55,8 +29,19 @@ void refCases(int AS) {
(void)p;
(void)s;
});
+}
+} // namespace badref1
+
+// Check reference kernel arguments witin structs or lambdas;
+namespace badref2 {
+struct S { // expected-note 2{{within field of type 'S' declared here}}
+ int a;
+ int &b; //expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
+};
- auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+void test() {
+ int p = 0;
+ auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
// expected-note at -1 {{within field of type}}
S Str {p, p};
kernel_single_task<class KN<2>>( // expected-note {{requested here}}
@@ -70,6 +55,22 @@ void refCases(int AS) {
(void)Str;
});
+}
+} // namespace badref2
+
+// Check references within array kernel arguments.
+namespace badref3 {
+struct S { // expected-note {{within field of type 'S' declared here}}
+ int a;
+ int &b; //expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+};
+
+void fooarr(int (&arr)[5]) {
+}
+
+void test(int AS) {
+ int p = 0;
+ S Str {p, p};
S arr[2] = {Str, Str};
kernel_single_task<class KN<4>>( // expected-note {{requested here}}
[=] { // expected-note {{within field of type}}
@@ -85,9 +86,37 @@ void refCases(int AS) {
[&] { // expected-note {{within field of type}}
fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
});
+}
+} // namespace badref3
+
+// Check callable objects containing references.
+namespace badref4 {
+template <typename T> class Callable { // expected-note 2{{within field of type 'Callable<int &>' declared here}}
+ T data; // expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
+public:
+ Callable(T d) : data(d) {}
+ void operator()() {
+ }
+};
+
+class Derived1 : Callable<int> { // expected-note {{within field of type 'Derived1' declared here}}
+ int &a; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+public:
+ Derived1(int d, int &b) : Callable<int>(d), a(b) {}
+};
+
+class Derived2 : Callable<int&> { // expected-note {{within base of type 'Callable<int &>' declared here}}
+ int a;
+public:
+ Derived2(int d, int &b) : Callable<int&>(b), a(d) {}
+};
+
+void test(int AS) {
+ int p = 0;
kernel_single_task<class KN<8>>(Callable<int&>{p}); // expected-note {{requested here}}
kernel_single_task<class KN<9>>(Callable<int>{p});
kernel_single_task<class KN<10>>(Derived1{p, p}); // expected-note {{requested here}}
kernel_single_task<class KN<11>>(Derived2{p, p}); // expected-note {{requested here}}
}
+} // namespace badref4
>From e2394464623225cec70d8684f85ddacd8997e5ec Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Wed, 27 May 2026 07:53:08 -0700
Subject: [PATCH 15/39] Make format happy
---
clang/include/clang/AST/SubobjectVisitor.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index ffd14b95d7dec..fdf1c5c681356 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -51,7 +51,7 @@ class SubobjectVisitorBase {
void traverseRecord(ptr_t<RecordDecl> RD) {
if (ptr_t<CXXRecordDecl> CRD = dyn_cast<CXXRecordDecl>(RD)) {
- for (non_ptr_t<CXXBaseSpecifier&> BS : CRD->bases()) {
+ for (non_ptr_t<CXXBaseSpecifier &> BS : CRD->bases()) {
if (getDerived().visitBaseSpecifierPre(&BS))
getDerived().visit(BS.getType());
getDerived().visitBaseSpecifierPost(&BS);
>From c3d8438e5f43d454835753a1026cf170970ee433 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 28 May 2026 07:02:07 -0700
Subject: [PATCH 16/39] Arguments -> parameters
---
clang/lib/Sema/SemaSYCL.cpp | 16 ++++++++--------
...ns.cpp => sycl-kernel-param-restrictions.cpp} | 4 ++--
2 files changed, 10 insertions(+), 10 deletions(-)
rename clang/test/SemaSYCL/{sycl-kernel-arg-restrictions.cpp => sycl-kernel-param-restrictions.cpp} (97%)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index a39c29cd45798..5e966bc7103d7 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -666,7 +666,7 @@ OutlinedFunctionDecl *BuildSYCLKernelEntryPointOutline(Sema &SemaRef,
return OFD;
}
-class KernelArgsChecker : public ConstSubobjectVisitor<KernelArgsChecker> {
+class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SemaSYCL &SemaSYCLRef;
bool IsValid = true;
using ObjectAccess =
@@ -694,8 +694,8 @@ class KernelArgsChecker : public ConstSubobjectVisitor<KernelArgsChecker> {
}
public:
- KernelArgsChecker(SemaSYCL &SR, SourceLocation Loc)
- : ConstSubobjectVisitor<KernelArgsChecker>(SR.getASTContext()),
+ KernelParamsChecker(SemaSYCL &SR, SourceLocation Loc)
+ : ConstSubobjectVisitor<KernelParamsChecker>(SR.getASTContext()),
SemaSYCLRef(SR) {}
void checkParameter(const ParmVarDecl *PVD) {
@@ -726,8 +726,8 @@ class KernelArgsChecker : public ConstSubobjectVisitor<KernelArgsChecker> {
auto DirectParent = ObjectAccessPath.back();
// Reference cannot be a base, so just assume we came via a FieldDecl.
if (isa<const ParmVarDecl *>(DirectParent)) {
- // If reference is a kernel argument, there is nothing to do. We allow
- // references in direct kernel arguments for better performance of the
+ // If reference is a kernel parameter, there is nothing to do. We allow
+ // references in direct kernel parameters for better performance of the
// host code and we eliminate them when building actual kernel.
return true;
}
@@ -752,8 +752,8 @@ class KernelArgsChecker : public ConstSubobjectVisitor<KernelArgsChecker> {
bool isInvalid() { return !IsValid; }
};
-bool verifyKernelArguments(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
- KernelArgsChecker KAC(SemaSYCLRef, FD->getLocation());
+bool verifyKernelParams(FunctionDecl *FD, SemaSYCL &SemaSYCLRef) {
+ KernelParamsChecker KAC(SemaSYCLRef, FD->getLocation());
for (auto Param : FD->parameters())
KAC.checkParameter(Param);
return KAC.isInvalid();
@@ -784,7 +784,7 @@ StmtResult SemaSYCL::BuildSYCLKernelCallStmt(FunctionDecl *FD,
getASTContext().getSYCLKernelInfo(SKEPAttr->getKernelName());
assert(declaresSameEntity(SKI.getKernelEntryPointDecl(), FD) &&
"SYCL kernel name conflict");
- if (verifyKernelArguments(FD, *this))
+ if (verifyKernelParams(FD, *this))
return StmtError();
// Build the outline of the synthesized device entry point function.
diff --git a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
similarity index 97%
rename from clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
rename to clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 2f7681de9c74e..5d549af450f32 100644
--- a/clang/test/SemaSYCL/sycl-kernel-arg-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -32,7 +32,7 @@ void test() {
}
} // namespace badref1
-// Check reference kernel arguments witin structs or lambdas;
+// Check reference kernel parameters witin structs or lambdas;
namespace badref2 {
struct S { // expected-note 2{{within field of type 'S' declared here}}
int a;
@@ -58,7 +58,7 @@ void test() {
}
} // namespace badref2
-// Check references within array kernel arguments.
+// Check references within array kernel parameters.
namespace badref3 {
struct S { // expected-note {{within field of type 'S' declared here}}
int a;
>From 515a864f137f4884cb265bd2570ffd153689736e Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
Date: Fri, 29 May 2026 13:10:10 +0200
Subject: [PATCH 17/39] Apply suggestions from code review
Co-authored-by: Tom Honermann <tom at honermann.net>
---
clang/include/clang/AST/SubobjectVisitor.h | 9 +++++++++
clang/lib/Sema/SemaSYCL.cpp | 4 ++++
clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp | 2 +-
3 files changed, 14 insertions(+), 1 deletion(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index fdf1c5c681356..92950efc370f6 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -18,6 +18,15 @@
namespace clang {
+/// Given a type, subobject visitors visit all subobjects of the type in depth
+/// first order. Both pre-order and post-order visitation are performed so that
+/// derived classes can maintain an access path to the visited elements.
+/// Subobjects include all base classes and non-static data members, including
+/// those that are not subobjects according to the C++standard like data
+/// members with a reference type. Virtual base classes are visited each time
+/// they appear in a class hierarchy despite there being only one actual
+/// subobject present in an object of a most derived type. Array elements are
+/// not individually visited; only their containing array is.
template <template <typename> class Ptr, typename Derived>
class SubobjectVisitorBase {
ASTContext &Ctx;
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 5e966bc7103d7..3f0613d088cfd 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -738,6 +738,10 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
<< DirectFieldParent->getType();
emitObjectAccessPathNotes();
+ // Don't visit the type of the reference since any further invalid
+ // kernel parameter types contained within the referenced type
+ // might not be relevant once the programmer addresses the
+ // invalid use of a reference.
IsValid = false;
return false;
}
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 5d549af450f32..e4d92a82c635d 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -47,7 +47,7 @@ void test() {
kernel_single_task<class KN<2>>( // expected-note {{requested here}}
[=] { // expected-note 2{{within field of type}}
(void)L;
- (void)Str; // no error because fail for L already
+ (void)Str;
});
kernel_single_task<class KN<3>>( // expected-note {{requested here}}
>From 9f868d5f487255597f776590475963505b83b301 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 29 May 2026 06:53:18 -0700
Subject: [PATCH 18/39] Add requested test cases.
---
.../sycl-kernel-param-restrictions.cpp | 39 +++++++++++++++++++
1 file changed, 39 insertions(+)
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index e4d92a82c635d..645e4e2884a5d 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -120,3 +120,42 @@ void test(int AS) {
}
} // namespace badref4
+
+// Test virtual bases.
+// FIXME: explicitly diagnose virtual bases within kernel parameters.
+namespace badref5 {
+
+class Base { // expected-note {{within field of type 'Base' declared here}}}
+ int &data; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
+public:
+ Base(int &a) : data(a) {}
+};
+
+class Derived : virtual Base { // expected-note {{within base of type 'Base' declared here}}
+public:
+ Derived(int &a) : Base(a) {}
+
+};
+
+void test() {
+ int p = 0;
+ kernel_single_task<class KN<12>>(Derived{p}); // expected-note {{requested here}}
+}
+} // namespace badref5
+
+// Check that a struct that hold a reference and captured by reference by lambda
+// kernel object is diagnosed correctly.
+namespace badref6 {
+void test() {
+ int a;
+ struct S {
+ int &dm;
+ };
+ S s {a};
+ kernel_single_task<class KN<13>>([&] { (void)s; });
+ // expected-error at -1 {{'S &' cannot be used as the type of a kernel parameter}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<13>, (lambda at {{.*}} requested here}}
+ // expected-note-re at -3 {{within field of type '(lambda at {{.*}} declared here}}
+}
+
+} // namespace badref7
>From 123a096e26a6444bc37bc5908a5112bc47f327b3 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 29 May 2026 07:06:33 -0700
Subject: [PATCH 19/39] Transform requested here notes
---
.../sycl-kernel-param-restrictions.cpp | 32 ++++++++++++-------
1 file changed, 21 insertions(+), 11 deletions(-)
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 645e4e2884a5d..3b5a45c72c40d 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -19,7 +19,8 @@ void test() {
int p = 0;
double q = 0;
float s = 0;
- kernel_single_task<class KN<1>>( // expected-note {{requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<1>>(
[ // expected-note2{{within field of type}}
// expected-error at +1 {{'int &' cannot be used as the type of a kernel parameter}}
&p, q,
@@ -44,13 +45,15 @@ void test() {
auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
// expected-note at -1 {{within field of type}}
S Str {p, p};
- kernel_single_task<class KN<2>>( // expected-note {{requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<2>>(
[=] { // expected-note 2{{within field of type}}
(void)L;
(void)Str;
});
- kernel_single_task<class KN<3>>( // expected-note {{requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<3>>(
[=] { // // expected-note {{within field of type}}
(void)Str;
});
@@ -72,17 +75,20 @@ void test(int AS) {
int p = 0;
S Str {p, p};
S arr[2] = {Str, Str};
- kernel_single_task<class KN<4>>( // expected-note {{requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<4>>(
[=] { // expected-note {{within field of type}}
(void)arr;
});
int arr1[AS];
- kernel_single_task<class KN<5>>( // expected-note {{requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<5>>(
[&] { // expected-note {{within field of type}}
(void)arr1; // expected-error {{'int (&)[AS]' cannot be used as the type of a kernel parameter}}
});
int arrayints[5] = {0};
- kernel_single_task<class KN<7>>( // expected-note {{requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<7>>(
[&] { // expected-note {{within field of type}}
fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
});
@@ -113,10 +119,13 @@ class Derived2 : Callable<int&> { // expected-note {{within base of type 'Callab
void test(int AS) {
int p = 0;
- kernel_single_task<class KN<8>>(Callable<int&>{p}); // expected-note {{requested here}}
+ kernel_single_task<class KN<8>>(Callable<int&>{p});
+ // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<9>>(Callable<int>{p});
- kernel_single_task<class KN<10>>(Derived1{p, p}); // expected-note {{requested here}}
- kernel_single_task<class KN<11>>(Derived2{p, p}); // expected-note {{requested here}}
+ kernel_single_task<class KN<10>>(Derived1{p, p});
+ // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ kernel_single_task<class KN<11>>(Derived2{p, p});
+ // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
}
} // namespace badref4
@@ -139,7 +148,8 @@ class Derived : virtual Base { // expected-note {{within base of type 'Base' dec
void test() {
int p = 0;
- kernel_single_task<class KN<12>>(Derived{p}); // expected-note {{requested here}}
+ kernel_single_task<class KN<12>>(Derived{p});
+ // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
}
} // namespace badref5
@@ -154,7 +164,7 @@ void test() {
S s {a};
kernel_single_task<class KN<13>>([&] { (void)s; });
// expected-error at -1 {{'S &' cannot be used as the type of a kernel parameter}}
- // expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<13>, (lambda at {{.*}} requested here}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
// expected-note-re at -3 {{within field of type '(lambda at {{.*}} declared here}}
}
>From f848eb619c75c41fe7b03b56d728b7a86536e269 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 1 Jun 2026 06:05:26 -0700
Subject: [PATCH 20/39] Captures are not fields now
They are but we wont tell anyone
---
clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 ++
clang/lib/Sema/SemaSYCL.cpp | 11 ++++++++---
.../SemaSYCL/sycl-kernel-param-restrictions.cpp | 16 ++++++++--------
3 files changed, 18 insertions(+), 11 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index bc0eb3b6b6d73..14bdd12bdf99b 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11802,6 +11802,8 @@ def err_record_with_pointers_kernel_param : Error<
"%select{struct|union}0 kernel parameters may not contain pointers">;
def note_within_field_of_type : Note<
"within field of type %0 declared here">;
+def note_within_capture : Note<
+ "within captured variable of lambda declared here">;
def note_within_base_of_type : Note<
"within base of type %0 declared here">;
def note_illegal_field_declared_here : Note<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 3f0613d088cfd..6aeba40e4d5c2 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -677,9 +677,14 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
void emitObjectAccessPathNotes() {
for (auto Parent : ObjectAccessPath) {
if (auto *FD = Parent.dyn_cast<const FieldDecl *>()) {
- SemaSYCLRef.Diag(FD->getParent()->getLocation(),
- diag::note_within_field_of_type)
- << FD->getParent();
+ if (FD->getParent()->isLambda()) {
+ SemaSYCLRef.Diag(FD->getParent()->getLocation(),
+ diag::note_within_capture);
+ } else {
+ SemaSYCLRef.Diag(FD->getParent()->getLocation(),
+ diag::note_within_field_of_type)
+ << FD->getParent();
+ }
} else if (auto *BS = Parent.dyn_cast<const CXXBaseSpecifier *>()) {
CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
assert(RD);
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 3b5a45c72c40d..1d2598b9d7c93 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -21,7 +21,7 @@ void test() {
float s = 0;
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<1>>(
- [ // expected-note2{{within field of type}}
+ [ // expected-note2{{within captured variable of lambda declared here}}
// expected-error at +1 {{'int &' cannot be used as the type of a kernel parameter}}
&p, q,
// expected-error at +1 {{'float &' cannot be used as the type of a kernel parameter}}
@@ -43,18 +43,18 @@ struct S { // expected-note 2{{within field of type 'S' declared here}}
void test() {
int p = 0;
auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
- // expected-note at -1 {{within field of type}}
+ // expected-note at -1 {{within captured variable of lambda declared here}}
S Str {p, p};
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<2>>(
- [=] { // expected-note 2{{within field of type}}
+ [=] { // expected-note 2{{within captured variable of lambda declared here}}
(void)L;
(void)Str;
});
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<3>>(
- [=] { // // expected-note {{within field of type}}
+ [=] { // // expected-note {{within captured variable of lambda declared here}}
(void)Str;
});
@@ -77,19 +77,19 @@ void test(int AS) {
S arr[2] = {Str, Str};
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<4>>(
- [=] { // expected-note {{within field of type}}
+ [=] { // expected-note {{within captured variable of lambda declared here}}
(void)arr;
});
int arr1[AS];
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<5>>(
- [&] { // expected-note {{within field of type}}
+ [&] { // expected-note {{within captured variable of lambda declared here}}
(void)arr1; // expected-error {{'int (&)[AS]' cannot be used as the type of a kernel parameter}}
});
int arrayints[5] = {0};
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<7>>(
- [&] { // expected-note {{within field of type}}
+ [&] { // expected-note {{within captured variable of lambda declared here}}
fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
});
}
@@ -165,7 +165,7 @@ void test() {
kernel_single_task<class KN<13>>([&] { (void)s; });
// expected-error at -1 {{'S &' cannot be used as the type of a kernel parameter}}
// expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
- // expected-note-re at -3 {{within field of type '(lambda at {{.*}} declared here}}
+ // expected-note at -3 {{within captured variable of lambda declared here}}
}
} // namespace badref7
>From f025ea78b73d398324600bd8047d5c733668d8bc Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 1 Jun 2026 06:10:52 -0700
Subject: [PATCH 21/39] Use auto instead of non_ptr_t
---
clang/include/clang/AST/SubobjectVisitor.h | 4 +---
1 file changed, 1 insertion(+), 3 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 92950efc370f6..9eeaff4a55735 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -31,8 +31,6 @@ template <template <typename> class Ptr, typename Derived>
class SubobjectVisitorBase {
ASTContext &Ctx;
template <typename Class> using ptr_t = typename Ptr<Class>::type;
- template <typename Class>
- using non_ptr_t = typename std::remove_pointer<ptr_t<Class>>::type;
public:
SubobjectVisitorBase(ASTContext &Ctx) : Ctx(Ctx) {}
@@ -60,7 +58,7 @@ class SubobjectVisitorBase {
void traverseRecord(ptr_t<RecordDecl> RD) {
if (ptr_t<CXXRecordDecl> CRD = dyn_cast<CXXRecordDecl>(RD)) {
- for (non_ptr_t<CXXBaseSpecifier &> BS : CRD->bases()) {
+ for (auto &BS : CRD->bases()) {
if (getDerived().visitBaseSpecifierPre(&BS))
getDerived().visit(BS.getType());
getDerived().visitBaseSpecifierPost(&BS);
>From 76f9e65591b5503bb2fb4c535616cb3d6229d84b Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 2 Jun 2026 06:48:32 -0700
Subject: [PATCH 22/39] Correct within lambda diagnostic
---
.../clang/Basic/DiagnosticSemaKinds.td | 2 +-
clang/lib/Sema/SemaSYCL.cpp | 11 +++----
.../sycl-kernel-param-restrictions.cpp | 29 ++++++++++++++-----
3 files changed, 28 insertions(+), 14 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 158149ca3156c..abb374040a371 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11803,7 +11803,7 @@ def err_record_with_pointers_kernel_param : Error<
def note_within_field_of_type : Note<
"within field of type %0 declared here">;
def note_within_capture : Note<
- "within captured variable of lambda declared here">;
+ "within capture %0 of lambda expression here">;
def note_within_base_of_type : Note<
"within base of type %0 declared here">;
def note_illegal_field_declared_here : Note<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 6aeba40e4d5c2..2a2d3f4625c3e 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -677,13 +677,14 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
void emitObjectAccessPathNotes() {
for (auto Parent : ObjectAccessPath) {
if (auto *FD = Parent.dyn_cast<const FieldDecl *>()) {
- if (FD->getParent()->isLambda()) {
- SemaSYCLRef.Diag(FD->getParent()->getLocation(),
- diag::note_within_capture);
+ const CXXRecordDecl *ParentRD = cast<CXXRecordDecl>(FD->getParent());
+ if (ParentRD->isLambda()) {
+ SemaSYCLRef.Diag(ParentRD->getLocation(), diag::note_within_capture)
+ << ParentRD->getCapture(FD->getFieldIndex())->getCapturedVar();
} else {
- SemaSYCLRef.Diag(FD->getParent()->getLocation(),
+ SemaSYCLRef.Diag(ParentRD->getLocation(),
diag::note_within_field_of_type)
- << FD->getParent();
+ << ParentRD;
}
} else if (auto *BS = Parent.dyn_cast<const CXXBaseSpecifier *>()) {
CXXRecordDecl *RD = BS->getType()->getAsCXXRecordDecl();
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 1d2598b9d7c93..479a5f19fffa4 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -21,7 +21,8 @@ void test() {
float s = 0;
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<1>>(
- [ // expected-note2{{within captured variable of lambda declared here}}
+ [ // expected-note{{within capture 'p' of lambda expression here}}
+ // expected-note at -1{{within capture 's' of lambda expression here}}
// expected-error at +1 {{'int &' cannot be used as the type of a kernel parameter}}
&p, q,
// expected-error at +1 {{'float &' cannot be used as the type of a kernel parameter}}
@@ -43,18 +44,19 @@ struct S { // expected-note 2{{within field of type 'S' declared here}}
void test() {
int p = 0;
auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
- // expected-note at -1 {{within captured variable of lambda declared here}}
+ // expected-note at -1 {{within capture 'p' of lambda expression here}}
S Str {p, p};
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<2>>(
- [=] { // expected-note 2{{within captured variable of lambda declared here}}
+ [=] { // expected-note {{within capture 'L' of lambda expression here}}
+ // expected-note at -1 {{within capture 'Str' of lambda expression here}}
(void)L;
(void)Str;
});
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<3>>(
- [=] { // // expected-note {{within captured variable of lambda declared here}}
+ [=] { // // expected-note {{within capture 'Str' of lambda expression here}}
(void)Str;
});
@@ -77,19 +79,19 @@ void test(int AS) {
S arr[2] = {Str, Str};
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<4>>(
- [=] { // expected-note {{within captured variable of lambda declared here}}
+ [=] { // expected-note {{within capture 'arr' of lambda expression here}}
(void)arr;
});
int arr1[AS];
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<5>>(
- [&] { // expected-note {{within captured variable of lambda declared here}}
+ [&] { // expected-note {{within capture 'arr1' of lambda expression here}}
(void)arr1; // expected-error {{'int (&)[AS]' cannot be used as the type of a kernel parameter}}
});
int arrayints[5] = {0};
// expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<7>>(
- [&] { // expected-note {{within captured variable of lambda declared here}}
+ [&] { // expected-note {{within capture 'arrayints' of lambda expression here}}
fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
});
}
@@ -165,7 +167,18 @@ void test() {
kernel_single_task<class KN<13>>([&] { (void)s; });
// expected-error at -1 {{'S &' cannot be used as the type of a kernel parameter}}
// expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
- // expected-note at -3 {{within captured variable of lambda declared here}}
+ // expected-note at -3 {{within capture 's' of lambda expression here}}
+}
+} // namespace badref6
+
+// Check that init capture is diagnosed correctly.
+namespace badref7 {
+void test() {
+ int p = 0;
+ kernel_single_task<class KN<14>>([&x=p] { (void)x; });
+ // expected-error at -1 {{'int &' cannot be used as the type of a kernel parameter}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note at -3 {{within capture 'x' of lambda expression here}}
}
} // namespace badref7
>From 9d2d9c33130b56b8dd312be9bec27b246f676e81 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 2 Jun 2026 10:01:25 -0700
Subject: [PATCH 23/39] Add ParmVarDecl note
---
.../clang/Basic/DiagnosticSemaKinds.td | 2 +
clang/lib/Sema/SemaSYCL.cpp | 6 +-
.../sycl-kernel-param-restrictions.cpp | 61 ++++++++++++++-----
3 files changed, 51 insertions(+), 18 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index abb374040a371..8a155c0a34599 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11806,6 +11806,8 @@ def note_within_capture : Note<
"within capture %0 of lambda expression here">;
def note_within_base_of_type : Note<
"within base of type %0 declared here">;
+def note_within_param_of_type : Note<
+ "within parameter %0 of type %1 declared here">;
def note_illegal_field_declared_here : Note<
"field of illegal %select{type|pointer type}0 %1 declared here">;
def err_opencl_type_struct_or_union_field : Error<
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 2a2d3f4625c3e..30ffecf868af9 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -692,9 +692,9 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SemaSYCLRef.Diag(BS->getBeginLoc(), diag::note_within_base_of_type)
<< RD;
} else {
- // Nothing to emit for ParmVarDecl since its location just points to
- // skep-attributed function template.
- assert(isa<const ParmVarDecl *>(Parent));
+ auto *Param = cast<const ParmVarDecl *>(Parent);
+ SemaSYCLRef.Diag(Param->getBeginLoc(), diag::note_within_param_of_type)
+ << Param << Param->getType();
}
}
}
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 479a5f19fffa4..5f4e9bf1905cb 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -8,18 +8,18 @@ template<int, int = 0> struct KN;
template<typename KNT, typename... Ts>
void sycl_kernel_launch(const char *, Ts...) {}
+// Check that reference captures of kernel that defined as lambda are diagnosed.
+namespace badref1 {
// Kernel entry point template definition.
template<typename KNT, typename T>
[[clang::sycl_kernel_entry_point(KNT)]]
-void kernel_single_task(T) {}
+void kernel_single_task(T t) {} // expected-note-re 2{{within parameter 't' of type '(lambda at {{.*}})' declared here}}
-// Check that reference captures of kernel that defined as lambda are diagnosed.
-namespace badref1 {
void test() {
int p = 0;
double q = 0;
float s = 0;
- // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'badref1::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<1>>(
[ // expected-note{{within capture 'p' of lambda expression here}}
// expected-note at -1{{within capture 's' of lambda expression here}}
@@ -36,6 +36,11 @@ void test() {
// Check reference kernel parameters witin structs or lambdas;
namespace badref2 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} // expected-note-re 3{{within parameter 't' of type '(lambda at {{.*}})' declared here}}
+
struct S { // expected-note 2{{within field of type 'S' declared here}}
int a;
int &b; //expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
@@ -46,7 +51,7 @@ void test() {
auto L = [&]() { (void)p;}; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
// expected-note at -1 {{within capture 'p' of lambda expression here}}
S Str {p, p};
- // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'badref2::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<2>>(
[=] { // expected-note {{within capture 'L' of lambda expression here}}
// expected-note at -1 {{within capture 'Str' of lambda expression here}}
@@ -54,7 +59,7 @@ void test() {
(void)Str;
});
- // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'badref2::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<3>>(
[=] { // // expected-note {{within capture 'Str' of lambda expression here}}
(void)Str;
@@ -65,6 +70,11 @@ void test() {
// Check references within array kernel parameters.
namespace badref3 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} // expected-note-re 3{{within parameter 't' of type '(lambda at {{.*}})' declared here}}
+
struct S { // expected-note {{within field of type 'S' declared here}}
int a;
int &b; //expected-error {{'int &' cannot be used as the type of a kernel parameter}}
@@ -77,19 +87,19 @@ void test(int AS) {
int p = 0;
S Str {p, p};
S arr[2] = {Str, Str};
- // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'badref3::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<4>>(
[=] { // expected-note {{within capture 'arr' of lambda expression here}}
(void)arr;
});
int arr1[AS];
- // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'badref3::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<5>>(
[&] { // expected-note {{within capture 'arr1' of lambda expression here}}
(void)arr1; // expected-error {{'int (&)[AS]' cannot be used as the type of a kernel parameter}}
});
int arrayints[5] = {0};
- // expected-note-re at +1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at +1 {{in instantiation of function template specialization 'badref3::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<7>>(
[&] { // expected-note {{within capture 'arrayints' of lambda expression here}}
fooarr(arrayints); // expected-error {{'int (&)[5]' cannot be used as the type of a kernel parameter}}
@@ -99,6 +109,13 @@ void test(int AS) {
// Check callable objects containing references.
namespace badref4 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} // expected-note {{within parameter 't' of type 'badref4::Callable<int &>' declared here}}
+ // expected-note at -1 {{within parameter 't' of type 'badref4::Derived1' declared here}}
+ // expected-note at -2 {{within parameter 't' of type 'badref4::Derived2' declared here}}
+
template <typename T> class Callable { // expected-note 2{{within field of type 'Callable<int &>' declared here}}
T data; // expected-error 2{{'int &' cannot be used as the type of a kernel parameter}}
public:
@@ -122,12 +139,12 @@ class Derived2 : Callable<int&> { // expected-note {{within base of type 'Callab
void test(int AS) {
int p = 0;
kernel_single_task<class KN<8>>(Callable<int&>{p});
- // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at -1 {{in instantiation of function template specialization 'badref4::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<9>>(Callable<int>{p});
kernel_single_task<class KN<10>>(Derived1{p, p});
- // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at -1 {{in instantiation of function template specialization 'badref4::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
kernel_single_task<class KN<11>>(Derived2{p, p});
- // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at -1 {{in instantiation of function template specialization 'badref4::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
}
} // namespace badref4
@@ -135,6 +152,10 @@ void test(int AS) {
// Test virtual bases.
// FIXME: explicitly diagnose virtual bases within kernel parameters.
namespace badref5 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} //expected-note {{within parameter 't' of type 'badref5::Derived' declared here}}
class Base { // expected-note {{within field of type 'Base' declared here}}}
int &data; // expected-error {{'int &' cannot be used as the type of a kernel parameter}}
@@ -151,13 +172,18 @@ class Derived : virtual Base { // expected-note {{within base of type 'Base' dec
void test() {
int p = 0;
kernel_single_task<class KN<12>>(Derived{p});
- // expected-note-re at -1 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at -1 {{in instantiation of function template specialization 'badref5::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
}
} // namespace badref5
// Check that a struct that hold a reference and captured by reference by lambda
// kernel object is diagnosed correctly.
namespace badref6 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} //expected-note-re {{within parameter 't' of type '(lambda at {{.*}})' declared here}}
+
void test() {
int a;
struct S {
@@ -166,18 +192,23 @@ void test() {
S s {a};
kernel_single_task<class KN<13>>([&] { (void)s; });
// expected-error at -1 {{'S &' cannot be used as the type of a kernel parameter}}
- // expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'badref6::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
// expected-note at -3 {{within capture 's' of lambda expression here}}
}
} // namespace badref6
// Check that init capture is diagnosed correctly.
namespace badref7 {
+// Kernel entry point template definition.
+template<typename KNT, typename T>
+[[clang::sycl_kernel_entry_point(KNT)]]
+void kernel_single_task(T t) {} //expected-note-re {{within parameter 't' of type '(lambda at {{.*}})' declared here}}
+
void test() {
int p = 0;
kernel_single_task<class KN<14>>([&x=p] { (void)x; });
// expected-error at -1 {{'int &' cannot be used as the type of a kernel parameter}}
- // expected-note-re at -2 {{in instantiation of function template specialization 'kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
+ // expected-note-re at -2 {{in instantiation of function template specialization 'badref7::kernel_single_task<KN<{{[0-9]+}}>, {{.*}}>' requested here}}
// expected-note at -3 {{within capture 'x' of lambda expression here}}
}
>From e7945ef4c60af00f71fce0a509baf35423c700f8 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 8 Jun 2026 08:33:17 -0700
Subject: [PATCH 24/39] Emit notes in reverse order
---
clang/lib/Sema/SemaSYCL.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp
index 30ffecf868af9..b942f19761f40 100644
--- a/clang/lib/Sema/SemaSYCL.cpp
+++ b/clang/lib/Sema/SemaSYCL.cpp
@@ -675,7 +675,7 @@ class KernelParamsChecker : public ConstSubobjectVisitor<KernelParamsChecker> {
SmallVector<ObjectAccess, 4> ObjectAccessPath;
void emitObjectAccessPathNotes() {
- for (auto Parent : ObjectAccessPath) {
+ for (auto Parent : llvm::reverse(ObjectAccessPath)) {
if (auto *FD = Parent.dyn_cast<const FieldDecl *>()) {
const CXXRecordDecl *ParentRD = cast<CXXRecordDecl>(FD->getParent());
if (ParentRD->isLambda()) {
>From 0587689ae6a9a101472fbd14fc9e9fa15adbdceb Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 8 Jun 2026 08:47:14 -0700
Subject: [PATCH 25/39] Base class
---
clang/include/clang/Basic/DiagnosticSemaKinds.td | 2 +-
clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp | 4 ++--
2 files changed, 3 insertions(+), 3 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index ad24ce8031674..e828355e3d1cc 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -11811,7 +11811,7 @@ def note_within_field_of_type : Note<
def note_within_capture : Note<
"within capture %0 of lambda expression here">;
def note_within_base_of_type : Note<
- "within base of type %0 declared here">;
+ "within base class of type %0 declared here">;
def note_within_param_of_type : Note<
"within parameter %0 of type %1 declared here">;
def note_illegal_field_declared_here : Note<
diff --git a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
index 5f4e9bf1905cb..66aa00da18a04 100644
--- a/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
+++ b/clang/test/SemaSYCL/sycl-kernel-param-restrictions.cpp
@@ -130,7 +130,7 @@ class Derived1 : Callable<int> { // expected-note {{within field of type 'Derive
Derived1(int d, int &b) : Callable<int>(d), a(b) {}
};
-class Derived2 : Callable<int&> { // expected-note {{within base of type 'Callable<int &>' declared here}}
+class Derived2 : Callable<int&> { // expected-note {{within base class of type 'Callable<int &>' declared here}}
int a;
public:
Derived2(int d, int &b) : Callable<int&>(b), a(d) {}
@@ -163,7 +163,7 @@ class Base { // expected-note {{within field of type 'Base' declared here}}}
Base(int &a) : data(a) {}
};
-class Derived : virtual Base { // expected-note {{within base of type 'Base' declared here}}
+class Derived : virtual Base { // expected-note {{within base class of type 'Base' declared here}}
public:
Derived(int &a) : Base(a) {}
>From 10825a8a22abfaea628886808eee7a22cf9f119f Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Thu, 11 Jun 2026 02:24:39 -0700
Subject: [PATCH 26/39] Add an assert for dependent types
---
clang/include/clang/AST/SubobjectVisitor.h | 2 ++
1 file changed, 2 insertions(+)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 9eeaff4a55735..6941897ed1efa 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -39,6 +39,8 @@ class SubobjectVisitorBase {
Derived &getDerived() { return *static_cast<Derived *>(this); }
void visit(QualType QT) {
+ assert(!QT->isDependentType());
+
// If the type is an array, visit its element type. Separate traversal of
// arrays is not needed because the array will be encountered as a
// FieldDecl.
>From fcef47399d75b9e3bd5a4a8e341e964015efda07 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Fri, 12 Jun 2026 07:12:47 -0700
Subject: [PATCH 27/39] Add unittest, handle atomic types and incomplete array
types
---
clang/include/clang/AST/SubobjectVisitor.h | 8 +-
clang/unittests/AST/CMakeLists.txt | 1 +
clang/unittests/AST/SubobjectVisitorTest.cpp | 134 +++++++++++++++++++
3 files changed, 140 insertions(+), 3 deletions(-)
create mode 100644 clang/unittests/AST/SubobjectVisitorTest.cpp
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 6941897ed1efa..b7307b1459ded 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -40,14 +40,16 @@ class SubobjectVisitorBase {
void visit(QualType QT) {
assert(!QT->isDependentType());
+ QT = QT.getDesugaredType(Ctx);
+
+ if (auto *ResAtomicType = QT->getAs<AtomicType>())
+ QT = ResAtomicType->getValueType();
// If the type is an array, visit its element type. Separate traversal of
// arrays is not needed because the array will be encountered as a
// FieldDecl.
-
if (QT->isArrayType()) {
- QualType ElTy =
- cast<ConstantArrayType>(Ctx.getAsArrayType(QT))->getElementType();
+ QualType ElTy = Ctx.getAsArrayType(QT)->getElementType();
getDerived().visit(ElTy);
return;
}
diff --git a/clang/unittests/AST/CMakeLists.txt b/clang/unittests/AST/CMakeLists.txt
index 3a12a4a06a33a..81010e0469685 100644
--- a/clang/unittests/AST/CMakeLists.txt
+++ b/clang/unittests/AST/CMakeLists.txt
@@ -34,6 +34,7 @@ add_clang_unittest(ASTTests
SourceLocationTest.cpp
StmtPrinterTest.cpp
StructuralEquivalenceTest.cpp
+ SubobjectVisitorTest.cpp
TemplateNameTest.cpp
TypePrinterTest.cpp
UnresolvedSetTest.cpp
diff --git a/clang/unittests/AST/SubobjectVisitorTest.cpp b/clang/unittests/AST/SubobjectVisitorTest.cpp
new file mode 100644
index 0000000000000..560fa71316be1
--- /dev/null
+++ b/clang/unittests/AST/SubobjectVisitorTest.cpp
@@ -0,0 +1,134 @@
+//===- unittests/AST/SubobjectVisitorTest.cpp - Subobject Visitor tests --===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include "clang/AST/SubobjectVisitor.h"
+#include "clang/AST/ASTContext.h"
+#include "clang/AST/Decl.h"
+#include "clang/AST/DeclCXX.h"
+#include "clang/ASTMatchers/ASTMatchFinder.h"
+#include "clang/Tooling/Tooling.h"
+#include "gtest/gtest.h"
+#include <string>
+#include <vector>
+
+using namespace clang;
+using namespace clang::ast_matchers;
+using namespace clang::tooling;
+
+namespace {
+
+// Helper class to record visited subobjects
+class RecordingVisitor : public SubobjectVisitor<RecordingVisitor> {
+public:
+ std::vector<std::string> VisitedBases;
+ std::vector<std::string> VisitedFields;
+
+ RecordingVisitor(ASTContext &Ctx) : SubobjectVisitor<RecordingVisitor>(Ctx) {}
+
+ bool visitBaseSpecifierPre(CXXBaseSpecifier *BS) {
+ std::string Name = BS->getType()->getAsCXXRecordDecl()->getNameAsString();
+ VisitedBases.push_back(Name);
+ return true;
+ }
+
+ bool visitFieldDeclPre(FieldDecl *FD) {
+ std::string Name = FD->getNameAsString();
+ VisitedFields.push_back(Name);
+ return true;
+ }
+
+};
+
+// Helper function to get a CXXRecordDecl by name
+const CXXRecordDecl *getCXXRecordDecl(ASTUnit *AST, const std::string &Name) {
+ auto Result = match(cxxRecordDecl(hasName(Name)).bind("record"),
+ AST->getASTContext());
+ if (Result.empty())
+ return nullptr;
+ return Result[0].getNodeAs<CXXRecordDecl>("record");
+}
+
+TEST(SubobjectVisitorTest, Basic) {
+ auto AST = buildASTFromCode(R"cpp(
+ struct Base {
+ int BaseF;
+ };
+ struct Inner {
+ int InnerF;
+ };
+ struct S : public Base {
+ int MemberF;
+ Inner StructF;
+ };
+ )cpp");
+ ASSERT_TRUE(AST.get());
+
+ const CXXRecordDecl *RD = getCXXRecordDecl(AST.get(), "S");
+ ASSERT_TRUE(RD);
+
+ RecordingVisitor Visitor(AST->getASTContext());
+ Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
+ std::nullopt, RD, false));
+
+ EXPECT_EQ(Visitor.VisitedBases.size(), 1u);
+ EXPECT_EQ(Visitor.VisitedFields.size(), 4u);
+ EXPECT_EQ(Visitor.VisitedFields[0], "BaseF");
+ EXPECT_EQ(Visitor.VisitedFields[1], "MemberF");
+ EXPECT_EQ(Visitor.VisitedFields[2], "StructF");
+ EXPECT_EQ(Visitor.VisitedFields[3], "InnerF");
+}
+
+TEST(SubobjectVisitorTest, Atomic) {
+ auto AST = buildASTFromCode(R"cpp(
+ struct S {
+ int SField;
+ };
+ struct T {
+ _Atomic S InnerAtomic;
+ };
+ )cpp");
+ ASSERT_TRUE(AST.get());
+
+ const CXXRecordDecl *RD = getCXXRecordDecl(AST.get(), "T");
+ ASSERT_TRUE(RD);
+
+ RecordingVisitor Visitor(AST->getASTContext());
+ Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
+ std::nullopt, RD, false));
+
+ EXPECT_EQ(Visitor.VisitedBases.size(), 0);
+ EXPECT_EQ(Visitor.VisitedFields.size(), 2);
+ EXPECT_EQ(Visitor.VisitedFields[0], "InnerAtomic");
+ EXPECT_EQ(Visitor.VisitedFields[1], "SField");
+}
+
+TEST(SubobjectVisitorTest, FAM) {
+ auto AST = buildASTFromCode(R"cpp(
+ struct S {
+ int SField[];
+ };
+ struct T {
+ S Inner;
+ };
+ )cpp");
+ ASSERT_TRUE(AST.get());
+
+ const CXXRecordDecl *RD = getCXXRecordDecl(AST.get(), "T");
+ ASSERT_TRUE(RD);
+
+ RecordingVisitor Visitor(AST->getASTContext());
+ Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
+ std::nullopt, RD, false));
+
+ EXPECT_EQ(Visitor.VisitedBases.size(), 0);
+ EXPECT_EQ(Visitor.VisitedFields.size(), 2);
+ EXPECT_EQ(Visitor.VisitedFields[0], "Inner");
+ EXPECT_EQ(Visitor.VisitedFields[1], "SField");
+}
+
+} // namespace
>From b6362573e5c03ca6d941be76930007c4605a4c88 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 02:48:22 -0700
Subject: [PATCH 28/39] Fix the warning about forward-declared ASTContext
---
clang/include/clang/AST/SubobjectVisitor.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index b7307b1459ded..e3da0289247c3 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -14,7 +14,7 @@
#ifndef LLVM_CLANG_AST_SUBOBJECTVISITOR_H
#define LLVM_CLANG_AST_SUBOBJECTVISITOR_H
-#include "clang/AST/Type.h"
+#include "clang/AST/ASTContext.h"
namespace clang {
>From ff04a25a741e65087e2cf5f1a349880c043157df Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 12:59:35 +0200
Subject: [PATCH 29/39] Apply suggestions from code review
Co-authored-by: Tom Honermann <tom at honermann.net>
---
clang/include/clang/AST/SubobjectVisitor.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index e3da0289247c3..38621ad161142 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -22,7 +22,7 @@ namespace clang {
/// first order. Both pre-order and post-order visitation are performed so that
/// derived classes can maintain an access path to the visited elements.
/// Subobjects include all base classes and non-static data members, including
-/// those that are not subobjects according to the C++standard like data
+/// those that are not subobjects according to the C++ standard like data
/// members with a reference type. Virtual base classes are visited each time
/// they appear in a class hierarchy despite there being only one actual
/// subobject present in an object of a most derived type. Array elements are
>From 251f612f417000021d130efd1e58959f94a25826 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 04:10:15 -0700
Subject: [PATCH 30/39] Verify the base in the unit test
---
clang/unittests/AST/SubobjectVisitorTest.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/clang/unittests/AST/SubobjectVisitorTest.cpp b/clang/unittests/AST/SubobjectVisitorTest.cpp
index 560fa71316be1..75ef3cba83a58 100644
--- a/clang/unittests/AST/SubobjectVisitorTest.cpp
+++ b/clang/unittests/AST/SubobjectVisitorTest.cpp
@@ -73,9 +73,10 @@ TEST(SubobjectVisitorTest, Basic) {
RecordingVisitor Visitor(AST->getASTContext());
Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
- std::nullopt, RD, false));
+ std::nullopt, RD, false));
EXPECT_EQ(Visitor.VisitedBases.size(), 1u);
+ EXPECT_EQ(Visitor.VisitedBases[0], "Base");
EXPECT_EQ(Visitor.VisitedFields.size(), 4u);
EXPECT_EQ(Visitor.VisitedFields[0], "BaseF");
EXPECT_EQ(Visitor.VisitedFields[1], "MemberF");
>From 870988237757e52d4b4c59fb66d63a7e3de19d68 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 04:13:44 -0700
Subject: [PATCH 31/39] Fix format
---
clang/unittests/AST/SubobjectVisitorTest.cpp | 9 ++++-----
1 file changed, 4 insertions(+), 5 deletions(-)
diff --git a/clang/unittests/AST/SubobjectVisitorTest.cpp b/clang/unittests/AST/SubobjectVisitorTest.cpp
index 75ef3cba83a58..913d9bda2f636 100644
--- a/clang/unittests/AST/SubobjectVisitorTest.cpp
+++ b/clang/unittests/AST/SubobjectVisitorTest.cpp
@@ -41,13 +41,12 @@ class RecordingVisitor : public SubobjectVisitor<RecordingVisitor> {
VisitedFields.push_back(Name);
return true;
}
-
};
// Helper function to get a CXXRecordDecl by name
const CXXRecordDecl *getCXXRecordDecl(ASTUnit *AST, const std::string &Name) {
- auto Result = match(cxxRecordDecl(hasName(Name)).bind("record"),
- AST->getASTContext());
+ auto Result =
+ match(cxxRecordDecl(hasName(Name)).bind("record"), AST->getASTContext());
if (Result.empty())
return nullptr;
return Result[0].getNodeAs<CXXRecordDecl>("record");
@@ -100,7 +99,7 @@ TEST(SubobjectVisitorTest, Atomic) {
RecordingVisitor Visitor(AST->getASTContext());
Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
- std::nullopt, RD, false));
+ std::nullopt, RD, false));
EXPECT_EQ(Visitor.VisitedBases.size(), 0);
EXPECT_EQ(Visitor.VisitedFields.size(), 2);
@@ -124,7 +123,7 @@ TEST(SubobjectVisitorTest, FAM) {
RecordingVisitor Visitor(AST->getASTContext());
Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
- std::nullopt, RD, false));
+ std::nullopt, RD, false));
EXPECT_EQ(Visitor.VisitedBases.size(), 0);
EXPECT_EQ(Visitor.VisitedFields.size(), 2);
>From 875fbba305032a22aff343bb52c1588a3b5e46aa Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 04:14:20 -0700
Subject: [PATCH 32/39] Recurse for atomic value type
---
clang/include/clang/AST/SubobjectVisitor.h | 6 ++++--
1 file changed, 4 insertions(+), 2 deletions(-)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 38621ad161142..95c75ec3bf949 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -42,8 +42,10 @@ class SubobjectVisitorBase {
assert(!QT->isDependentType());
QT = QT.getDesugaredType(Ctx);
- if (auto *ResAtomicType = QT->getAs<AtomicType>())
- QT = ResAtomicType->getValueType();
+ if (auto *ResAtomicType = QT->getAs<AtomicType>()) {
+ getDerived().visit(ResAtomicType->getValueType());
+ return;
+ }
// If the type is an array, visit its element type. Separate traversal of
// arrays is not needed because the array will be encountered as a
>From 2aa1c33a47fea8d3627275032c7de8f189187fac Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 04:21:10 -0700
Subject: [PATCH 33/39] Test sugared atomic value type
---
clang/unittests/AST/SubobjectVisitorTest.cpp | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/clang/unittests/AST/SubobjectVisitorTest.cpp b/clang/unittests/AST/SubobjectVisitorTest.cpp
index 913d9bda2f636..b6e76f55f5a91 100644
--- a/clang/unittests/AST/SubobjectVisitorTest.cpp
+++ b/clang/unittests/AST/SubobjectVisitorTest.cpp
@@ -88,8 +88,9 @@ TEST(SubobjectVisitorTest, Atomic) {
struct S {
int SField;
};
+ typedef S STypDef;
struct T {
- _Atomic S InnerAtomic;
+ _Atomic STypDef InnerAtomic;
};
)cpp");
ASSERT_TRUE(AST.get());
>From a7d53513b13219a91838f1c3d97a4cef3c17e79b Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 13:24:22 +0200
Subject: [PATCH 34/39] Update clang/include/clang/AST/SubobjectVisitor.h
Co-authored-by: Tom Honermann <tom at honermann.net>
---
clang/include/clang/AST/SubobjectVisitor.h | 16 ++++++++++++++++
1 file changed, 16 insertions(+)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 95c75ec3bf949..1242c5b901bf1 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -42,6 +42,22 @@ class SubobjectVisitorBase {
assert(!QT->isDependentType());
QT = QT.getDesugaredType(Ctx);
+ // FunctionType, FunctionProtoType, and FunctionNoProtoType are never
+ // the type of a subobject.
+
+ // ObjCObjectType and ObjCInterfaceType are never the type of a
+ // subobject due to the Objective-C object allocation model.
+
+ // PointerType, BlockPointerType, ReferenceType, LValueReferenceType,
+ // RValueReferenceType, MemberPointerType, and ObjCObjectPointerType
+ // may all be the type of a subobject, but they do not contain other
+ // subobjects; the pointee type should not be visited.
+
+ // ComplexType, VectorType, ExtVectorType, MatrixType, PipeType,
+ // EnumType, and OverflowBehaviorType all have an element or
+ // underlying type that could be visited. However, in each of these
+ // cases, the lower type is constrained to a fundamental type and
+ // therefore doesn't contain any fields or base classes.
if (auto *ResAtomicType = QT->getAs<AtomicType>()) {
getDerived().visit(ResAtomicType->getValueType());
return;
>From 6fe7e11254b207a3033f99ad23774784f490dbc9 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Mon, 29 Jun 2026 05:41:39 -0700
Subject: [PATCH 35/39] Fix warning in the test
---
clang/unittests/AST/SubobjectVisitorTest.cpp | 8 ++++----
1 file changed, 4 insertions(+), 4 deletions(-)
diff --git a/clang/unittests/AST/SubobjectVisitorTest.cpp b/clang/unittests/AST/SubobjectVisitorTest.cpp
index b6e76f55f5a91..21c9cda579af6 100644
--- a/clang/unittests/AST/SubobjectVisitorTest.cpp
+++ b/clang/unittests/AST/SubobjectVisitorTest.cpp
@@ -102,8 +102,8 @@ TEST(SubobjectVisitorTest, Atomic) {
Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
std::nullopt, RD, false));
- EXPECT_EQ(Visitor.VisitedBases.size(), 0);
- EXPECT_EQ(Visitor.VisitedFields.size(), 2);
+ EXPECT_EQ(Visitor.VisitedBases.size(), 0u);
+ EXPECT_EQ(Visitor.VisitedFields.size(), 2u);
EXPECT_EQ(Visitor.VisitedFields[0], "InnerAtomic");
EXPECT_EQ(Visitor.VisitedFields[1], "SField");
}
@@ -126,8 +126,8 @@ TEST(SubobjectVisitorTest, FAM) {
Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
std::nullopt, RD, false));
- EXPECT_EQ(Visitor.VisitedBases.size(), 0);
- EXPECT_EQ(Visitor.VisitedFields.size(), 2);
+ EXPECT_EQ(Visitor.VisitedBases.size(), 0u);
+ EXPECT_EQ(Visitor.VisitedFields.size(), 2u);
EXPECT_EQ(Visitor.VisitedFields[0], "Inner");
EXPECT_EQ(Visitor.VisitedFields[1], "SField");
}
>From d3cf31c1dc2b80f4c20d28cf4f6c3c51dd8a7fb4 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva at intel.com>
Date: Tue, 30 Jun 2026 10:54:31 +0200
Subject: [PATCH 36/39] Update clang/include/clang/AST/SubobjectVisitor.h
Co-authored-by: Tom Honermann <tom at honermann.net>
---
clang/include/clang/AST/SubobjectVisitor.h | 1 +
1 file changed, 1 insertion(+)
diff --git a/clang/include/clang/AST/SubobjectVisitor.h b/clang/include/clang/AST/SubobjectVisitor.h
index 1242c5b901bf1..473446dd36b4b 100644
--- a/clang/include/clang/AST/SubobjectVisitor.h
+++ b/clang/include/clang/AST/SubobjectVisitor.h
@@ -58,6 +58,7 @@ class SubobjectVisitorBase {
// underlying type that could be visited. However, in each of these
// cases, the lower type is constrained to a fundamental type and
// therefore doesn't contain any fields or base classes.
+
if (auto *ResAtomicType = QT->getAs<AtomicType>()) {
getDerived().visit(ResAtomicType->getValueType());
return;
>From 922c3967136bc456203f9c8d4392b428d89e0f9d Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 30 Jun 2026 03:01:37 -0700
Subject: [PATCH 37/39] Update the test
---
clang/unittests/AST/SubobjectVisitorTest.cpp | 9 +++------
1 file changed, 3 insertions(+), 6 deletions(-)
diff --git a/clang/unittests/AST/SubobjectVisitorTest.cpp b/clang/unittests/AST/SubobjectVisitorTest.cpp
index 21c9cda579af6..5df382b976420 100644
--- a/clang/unittests/AST/SubobjectVisitorTest.cpp
+++ b/clang/unittests/AST/SubobjectVisitorTest.cpp
@@ -71,8 +71,7 @@ TEST(SubobjectVisitorTest, Basic) {
ASSERT_TRUE(RD);
RecordingVisitor Visitor(AST->getASTContext());
- Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
- std::nullopt, RD, false));
+ Visitor.visit(AST->getASTContext().getCanonicalTagType(RD));
EXPECT_EQ(Visitor.VisitedBases.size(), 1u);
EXPECT_EQ(Visitor.VisitedBases[0], "Base");
@@ -99,8 +98,7 @@ TEST(SubobjectVisitorTest, Atomic) {
ASSERT_TRUE(RD);
RecordingVisitor Visitor(AST->getASTContext());
- Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
- std::nullopt, RD, false));
+ Visitor.visit(AST->getASTContext().getCanonicalTagType(RD));
EXPECT_EQ(Visitor.VisitedBases.size(), 0u);
EXPECT_EQ(Visitor.VisitedFields.size(), 2u);
@@ -123,8 +121,7 @@ TEST(SubobjectVisitorTest, FAM) {
ASSERT_TRUE(RD);
RecordingVisitor Visitor(AST->getASTContext());
- Visitor.visit(AST->getASTContext().getTagType(ElaboratedTypeKeyword::None,
- std::nullopt, RD, false));
+ Visitor.visit(AST->getASTContext().getCanonicalTagType(RD));
EXPECT_EQ(Visitor.VisitedBases.size(), 0u);
EXPECT_EQ(Visitor.VisitedFields.size(), 2u);
>From 26557fcb86695807671332ff707f5dbdcb972404 Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 30 Jun 2026 06:03:15 -0700
Subject: [PATCH 38/39] Add a release note
---
clang/docs/ReleaseNotes.rst | 3 +++
1 file changed, 3 insertions(+)
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
index a2eef75e2a719..8c68fe8a59223 100644
--- a/clang/docs/ReleaseNotes.rst
+++ b/clang/docs/ReleaseNotes.rst
@@ -1183,6 +1183,9 @@ SYCL Support
library was unconditionally added to the link line even when the flag was
passed.
+- Clang now is capable of diagnosing reference kernel parameters which are not
+ allowed by SYCL 2020 spec.
+
Improvements
^^^^^^^^^^^^
- Improved substitution performance in concept checking. (#GH172266)
>From e2f3574ec9a8f74aa980da0e4778ac036819e38a Mon Sep 17 00:00:00 2001
From: "Podchishchaeva, Mariya" <mariya.podchishchaeva at intel.com>
Date: Tue, 30 Jun 2026 06:14:19 -0700
Subject: [PATCH 39/39] Remove the rst
---
clang/docs/ReleaseNotes.rst | 1209 -----------------------------------
1 file changed, 1209 deletions(-)
delete mode 100644 clang/docs/ReleaseNotes.rst
diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst
deleted file mode 100644
index 8c68fe8a59223..0000000000000
--- a/clang/docs/ReleaseNotes.rst
+++ /dev/null
@@ -1,1209 +0,0 @@
-.. If you want to modify sections/contents permanently, you should modify both
- ReleaseNotes.rst and ReleaseNotesTemplate.txt.
-
-===========================================
-Clang |release| |ReleaseNotesTitle|
-===========================================
-
-.. contents::
- :local:
- :depth: 2
-
-Written by the `LLVM Team <https://llvm.org/>`_
-
-.. only:: PreRelease
-
- .. warning::
- These are in-progress notes for the upcoming Clang |version| release.
- Release notes for previous releases can be found on
- `the Releases Page <https://llvm.org/releases/>`_.
-
-Introduction
-============
-
-This document contains the release notes for the Clang C/C++/Objective-C
-frontend, part of the LLVM Compiler Infrastructure, release |release|. Here we
-describe the status of Clang in some detail, including major
-improvements from the previous release and new feature work. For the
-general LLVM release notes, see `the LLVM
-documentation <https://llvm.org/docs/ReleaseNotes.html>`_. For the libc++ release notes,
-see `this page <https://libcxx.llvm.org/ReleaseNotes.html>`_. All LLVM releases
-may be downloaded from the `LLVM releases web site <https://llvm.org/releases/>`_.
-
-For more information about Clang or LLVM, including information about the
-latest release, please see the `Clang Web Site <https://clang.llvm.org>`_ or the
-`LLVM Web Site <https://llvm.org>`_.
-
-Potentially Breaking Changes
-============================
-
-C/C++ Language Potentially Breaking Changes
--------------------------------------------
-
-- Clang now makes it ill-formed to try to ``break`` out of or ``continue`` a loop inside its own condition,
- increment, or init-statement in all C and C++ language modes. This means that code such as
-
- .. code-block:: c++
-
- while (({ break; })) {}
-
- is now ill-formed. An outer loop can still be broken out of or continued so long as the inner loop is
- in the body of the outer loop:
-
- .. code-block:: c++
-
- // Ok, this breaks out of the 'for' loop.
- for (;;) {
- while (({ break; true; })) {}
- }
-
- // Error: can't break out of the 'for' loop from within its own increment.
- for (;;({ while (({ break; true; })) {} })) {}
-
- This also resolves a divergence from GCC: in a construct such as
-
- .. code-block:: c++
-
- for (;;) {
- while (({ break; true; })) {}
- }
-
- Clang would previously ``break`` out of the ``while`` loop, whereas GCC (since version 9) would
- ``break`` out of the ``for`` loop here. Now, Clang and GCC both break out of the ``for`` loop.
-
-- Clang now parses line and digit directives, module names, and original filenames as unevaluated
- strings. This means that code containing strings with escape sequences such as
-
- .. code-block:: c++
-
- # 1 "original\x12source.c"
- #pragma clang module import "\x41"
- # 50 "a\012.c"
-
- are now ill-formed.
-
-C++ Specific Potentially Breaking Changes
------------------------------------------
-
-- Clang now more aggressively optimizes away stores to objects after they are
- dead. This behavior can be disabled with ``-fno-lifetime-dse``.
-
-- Clang now correctly rejects ``export`` declarations in module implementation
- partitions. (#GH107602)
-
-- Template argument deduction now treats the ``N`` in ``_BitInt(N)``
- as being of type ``std::size_t`` instead of ``int``,
- matching the deduction of array sizes from ``int(&)[N]``.
- This is a breaking change for code that depended on the previously deduced type. (#GH195033)
-
-- Clang now rejects nested local classes defined in a different
- block scope than their parent class. (#GH193472)
-
-ABI Changes in This Version
----------------------------
-
-- Fixed incorrect struct layout for ``_BitInt`` bitfields wider than 255 bits
- on MSVC targets. Internal bitfield tracking fields were changed from
- ``unsigned char`` to ``uint64_t`` to prevent overflow. This might be an ABI
- break for such structs compared to earlier Clang versions.
-- Fixed a number of issues with the ``__regcall`` calling convention for passing
- structs on non-Windows x86-64 targets, including a crash when handling empty
- struct arguments. This changes how structs that contain arrays, floating point
- types, or ``_Complex float`` types are passed, and may introduce
- incompatibilities with code compiled by earlier versions of Clang that uses
- the ``__regcall`` calling convention on these targets. (#GH62999) (#GH98635)
-- Fixed Itanium mangling for lambdas in instantiated non-static data member
- initializers by preserving the field-name closure-prefix. This changes the
- mangled names for affected lambdas. (#GH190555)
-- Clang now uses MSVC-compatible manglings for supported AArch64 SVE builtin
- types when targeting the Microsoft ABI. This changes symbol names for
- affected declarations compared to previous Clang releases. (#GH196170)
-- The resume and destroy functions of C++20 coroutines now use the platform C
- calling convention instead of LLVM's internal ``fastcc``. This makes the
- coroutine ABI stable across LLVM versions and interoperable with other
- compilers. On most targets this is not a breaking change because ``fastcc``
- and the platform C calling convention agree for ``void(ptr)``. It is an ABI
- break on i686, MIPS O32, PowerPC64 ELFv1, and Lanai.
-- Fixed incorrect struct return when single large vector (256/512-bit) used on
- x86-64 targets. (#GH203760) The bug was introduced since Clang 21. (#GH120670)
-
-AST Dumping Potentially Breaking Changes
-----------------------------------------
-
-- The JSON AST dump now includes all fields from ``AvailabilityAttr``: ``platform``,
- ``introduced``, ``deprecated``, ``obsoleted``, ``unavailable``, ``message``,
- ``strict``, ``replacement``, ``priority``, and ``environment``. Previously, these
- fields were missing from the JSON output.
-- Colons that appear at the end of a ParamCommentCommand name are not serialized
- as part of the name.
-
-Clang Frontend Potentially Breaking Changes
--------------------------------------------
-
-- HIPSPV toolchain: `--offload-targets=spirv{32,64}` option is
- deprecated and will be removed when the new offload driver becomes
- default. The replacement for the option is
- `--offload-targets=spirv{32,64}-unknown-chipstar` when using the new
- offload driver (`--offload-new-driver`).
-- The new driver (`--offload-new-driver`) is now default for all offloading
- compilations. This changes the ABI for relocatable device code. Currently,
- libraries will need to be recompiled, or used with
- (`--no-offload-new-driver`). This option will be removed in the next release.
-
-- Clang no longer defines the ``__cpp_impl_coroutine`` feature test macro under the 32-bit x86 Microsoft ABI,
- as support for coroutines on this target is incomplete.
- When using coroutines on this target a warning is emmitted to indicate the lack of full support.
- That warning can be disabled with ``-Wno-coroutines-unsupported-target``. (see #GH59382)
-
-Clang Python Bindings Potentially Breaking Changes
---------------------------------------------------
-- Remove ``CompletionString.Availability``. No libclang interfaces returned instances of it.
-- ``CompletionString.availability`` now returns instances of ``CompletionString.AvailabilityKindCompat``.
-
- Instances of ``AvailabilityKindCompat`` have the same ``__str__`` representation
- as the previous ``CompletionChunk.Kind`` and are equality-comparable with
- the existing ``AvailabilityKind`` enum. It will be replaced by ``AvailabilityKind``
- in a future release. When this happens, the return type of ``CompletionString.availability``
- will change to ``AvailabilityKind``, so it is recommended to use ``AvailabilityKind``
- to compare with the return values of ``CompletionString.availability``.
-- Remove ``availabilityKinds``. In this release, uses of ``availabilityKinds``
- need to be replaced by ``CompletionString.AvailabilityKind``.
-- ``CompletionChunk.kind`` now returns instances of ``CompletionChunkKind``.
-
- Instances of ``CompletionChunkKind`` have the same ``__str__`` representation
- as the previous ``CompletionChunk.Kind`` for compatibility.
- These representations will be changed in a future release to match other enums.
-- Remove ``completionChunkKindMap``. In this release, uses of ``completionChunkKindMap``
- need to be replaced by ``CompletionChunkKind``.
-- Move ``SPELLING_CACHE`` into ``CompletionChunk`` and change it to use
- ``CompletionChunkKind`` instances as keys, instead of the enum values.
- An alias is kept in the form of a ``SPELLING_CACHE`` variable, but it only supports
- ``__getitem__`` and ``__contains__``. It will be removed in a future release.
- Please migrate to using ``CompletionChunk.SPELLING_CACHE`` instead.
-- ``SourceLocation`` and ``SourceRange`` now use ``NotImplemented`` to delegate
- equality checks (``__eq__``) to the other object they are compared with when
- they are of different classes. They previously returned ``False`` when compared
- with objects of other classes.
-- ``TranslationUnit.get_tokens`` now throws an error if both the ``extent`` and
- ``locations`` argument are passed. Previousy, ``locations`` took precedence.
-- ``_CXUnsavedFile`` will be renamed to ``UnsavedFile`` for consistency.
- ``UnsavedFile`` is already available to use and existing uses should
- be adapted to refer to it instead. ``_CXUnsavedFile`` will be removed in a
- future release.
-
-OpenCL Potentially Breaking Changes
------------------------------------
-- Clang now diagnoses zero-length arrays as errors in OpenCL. OpenCL C 3.0
- section 6.11.d states that "Variable length arrays and structures with
- flexible (or unsized) arrays are not supported."
-
-What's New in Clang |release|?
-==============================
-
-C++ Language Changes
---------------------
-
-- ``__is_trivially_equality_comparable`` no longer returns false for all enum types. (#GH132672)
-- ``auto`` parameters are now available in all C++ language modes as an extension.
-
-C++2c Feature Support
-^^^^^^^^^^^^^^^^^^^^^
-
-- Clang now propagates ``constinit`` and ``constexpr`` in structured bindings with tuple-like initializers.
-
-C++23 Feature Support
-^^^^^^^^^^^^^^^^^^^^^
-
-- Partially implement Itanium mangling for pack indexing. Partially substituted packs are not yet supported. (#GH112003)
-
-C++20 Feature Support
-^^^^^^^^^^^^^^^^^^^^^
-
-- Clang now supports `P1857R3 <https://wg21.link/p1857r3>`_ Modules Dependency Discovery. (#GH54047)
-
-C++17 Feature Support
-^^^^^^^^^^^^^^^^^^^^^
-
-Resolutions to C++ Defect Reports
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-- Implemented `CWG1780 Explicit instantiation/specialization of generic lambda
- operator() <https://cplusplus.github.io/CWG/issues/1780.html>`_
-
-- Clang now allows omitting ``typename`` before a template name in a
- conversion operator, implementing `CWG2413 <https://wg21.link/cwg2413>`_.
-
-- Clang now uses non-reference types for structured bindings whose initializer
- returns a prvalue. This resolves `CWG3135 <https://wg21.link/cwg3135>`_.
-
-C Language Changes
-------------------
-
-C2y Feature Support
-^^^^^^^^^^^^^^^^^^^
-
-- Implemented the type-specific C2y ``<stdbit.h>`` rotate functions with constexpr
- evaluation support:
- ``stdc_rotate_left_{uc,us,ui,ul,ull}`` and
- ``stdc_rotate_right_{uc,us,ui,ul,ull}``.
-
-- Implemented C2y ``<stdbit.h>`` memory reversal functions:
- ``__builtin_stdc_memreverse8`` / ``stdc_memreverse8`` (in-place byte
- reversal of a byte array) and ``stdc_memreverse8u{8,16,32,64}`` (byte-swap
- of an exact-width unsigned integer value, usable in constant expressions).
-
-C23 Feature Support
-^^^^^^^^^^^^^^^^^^^
-- Clang now allows C23 ``constexpr`` struct member access through the dot operator in constant expressions. (#GH178349)
-- Fixed a failing assertion when validating an invalid structure redefinition
- with a member which uses an incomplete enumeration type. (#GH190227)
-- Clang now supports the C23 ``wN`` and ``wfN`` length modifiers. (#GH116962)
-- Clang now recognizes the C23 ``H``, ``D``, and ``DD`` length modifiers in
- format strings and diagnoses their use because Clang does not yet support
- the corresponding decimal floating-point types, ``_Decimal32``, ``_Decimal64``, and ``_Decimal128``. (#GH116962)
-
-Objective-C Language Changes
------------------------------
-
-- Clang now emits Objective-C number, array, and dictionary literals as
- compile-time constant data structures rather than runtime ``objc_msgSend``
- calls on targets whose runtime supports constant literal classes. The
- feature can be disabled altogether with ``-fno-objc-constant-literals``,
- or selectively per literal kind with ``-fno-constant-nsnumber-literals``,
- ``-fno-constant-nsarray-literals``, and
- ``-fno-constant-nsdictionary-literals``.
-
-Non-comprehensive list of changes in this release
--------------------------------------------------
-
-- Added support for floating point and pointer values in most ``__atomic_``
- builtins.
-
-- Added ``__builtin_stdc_rotate_left`` and ``__builtin_stdc_rotate_right``
- for bit rotation of unsigned integers including ``_BitInt`` types. Rotation
- counts are normalized modulo the bit-width and support negative values.
- Usable in constant expressions. Implicit conversion is supported for
- class/struct types with conversion operators.
-- Implemented the following C23 ``<stdbit.h>`` builtins with ``_BitInt``
- support and constexpr evaluation:
- ``__builtin_stdc_leading_zeros``, ``__builtin_stdc_leading_ones``,
- ``__builtin_stdc_trailing_zeros``, ``__builtin_stdc_trailing_ones``,
- ``__builtin_stdc_first_leading_zero``, ``__builtin_stdc_first_leading_one``,
- ``__builtin_stdc_first_trailing_zero``, ``__builtin_stdc_first_trailing_one``,
- ``__builtin_stdc_count_zeros``, ``__builtin_stdc_count_ones``,
- ``__builtin_stdc_has_single_bit``, ``__builtin_stdc_bit_width``,
- ``__builtin_stdc_bit_floor``, and ``__builtin_stdc_bit_ceil``.
-
-- Implemented the type-specific C23 ``<stdbit.h>`` functions with constexpr
- evaluation support:
- ``stdc_leading_zeros_{uc,us,ui,ul,ull}``,
- ``stdc_leading_ones_{uc,us,ui,ul,ull}``,
- ``stdc_trailing_zeros_{uc,us,ui,ul,ull}``,
- ``stdc_trailing_ones_{uc,us,ui,ul,ull}``,
- ``stdc_first_leading_zero_{uc,us,ui,ul,ull}``,
- ``stdc_first_leading_one_{uc,us,ui,ul,ull}``,
- ``stdc_first_trailing_zero_{uc,us,ui,ul,ull}``,
- ``stdc_first_trailing_one_{uc,us,ui,ul,ull}``,
- ``stdc_count_zeros_{uc,us,ui,ul,ull}``,
- ``stdc_count_ones_{uc,us,ui,ul,ull}``,
- ``stdc_has_single_bit_{uc,us,ui,ul,ull}``,
- ``stdc_bit_width_{uc,us,ui,ul,ull}``,
- ``stdc_bit_floor_{uc,us,ui,ul,ull}``, and
- ``stdc_bit_ceil_{uc,us,ui,ul,ull}``.
-
-- A new generic bit-reverse builtin function ``__builtin_bitreverseg`` that
- extends bit-reversal support to all standard integers type, including
- ``_BitInt``
-
-- Added ``__builtin_elementwise_clmul`` for carry-less multiplication of
- integers including ``_BitInt`` types. This includes constexpr evaluation
- support.
-
-- Added ``__builtin_elementwise_pext`` and ``__builtin_elementwise_pdep`` for
- parallel bit extract and parallel bit deposit operations on integers including
- ``_BitInt`` types. This includes constexpr evaluation support.
-
-- Deprecated float types support from ``__builtin_elementwise_max`` and
- ``__builtin_elementwise_min``.
-
-- Added header ``endian.h`` which contains byte order helpers specified in POSIX
-
-- Added #pragma loop licm(disable) for llvm.loop.licm.disable metadata
-
-- Added a new ``ExplicitInstantiationDecl`` AST node to represent explicit
- template instantiations (e.g., ``template void foo<int>();`` or
- ``extern template class S<int>;``). Previously, source location information
- for explicit instantiation statements was discarded after parsing. The new
- node preserves the full source range including the ``extern`` and ``template``
- keywords, qualifiers, template arguments as written, and the declared type,
- enabling tools such as language servers and refactoring engines to accurately
- map source locations back to explicit instantiation sites.
-
-- ``typeid`` on references and pointers of ``final`` types no longer emits a
- vtable lookup at runtime.
-
-- Updated support for Unicode from 15.1 to 18.0.
-
-- Linux and Windows toolchains now support Clang multilibs using
- ``-fmultilib-flag=``.
-
-- The SafeStack builtins ``__builtin___get_unsafe_stack_ptr``,
- ``__builtin___get_unsafe_stack_bottom``, ``__builtin___get_unsafe_stack_top``,
- and ``__builtin___get_unsafe_stack_start`` are now deprecated. Use the
- corresponding functions from ``<sanitizer/safestack_interface.h>`` instead.
-
-New Compiler Flags
-------------------
-- New option ``-fms-anonymous-structs`` / ``-fno-ms-anonymous-structs`` added
- to enable or disable Microsoft's anonymous struct/union extension without
- enabling other ``-fms-extensions`` features (#GH177607).
-- New option ``--precompile-reduced-bmi`` allows build system to generate a
- reduced BMI only for a C++20 importable module unit. Previously the users
- can only generate the reduced BMI as a by-product, e.g, an object files or
- a full BMI.
-
-- New ``-cc1`` option ``-fexperimental-overflow-behavior-types`` added to
- enable parsing of the experimental ``overflow_behavior`` type attribute and
- type specifiers.
-
-- New ``-cl`` option ``/d2guardnochecks`` added to match MSVC. When Windows
- Control Flow Guard (CFG) is enabled by other options, it will instruct Clang
- to emit the CFG metadata, but disable adding checks.
-
-- New option ``-fdiagnostics-show-inlining-chain`` added to show inlining chain
- notes for ``[[gnu::warning]]`` and ``[[gnu::error]]`` diagnostics. When a
- function with these attributes is called from an inlined context, Clang can
- now show which functions were inlined to reach the call. When debug info is
- available (``-gline-directives-only`` (implicitly enabled at ``-g1``) or
- higher), accurate source locations are used; otherwise, a heuristic fallback
- is used with a note suggesting how to enable debug info for better accuracy.
-
-- New option ``-fwin-cfg-mechanism=`` added to control the mechanism used by
- Control Flow Guard on Windows. Accepted values are ``automatic`` (default),
- ``dispatch``, and ``check``. The ``dispatch`` mechanism uses the dispatch
- function to perform indirect call checks and can improve performance, while
- ``check`` uses the traditional check mechanism.
-- New ``-cl`` option ``/d2guardcfgdispatch`` added to match MSVC. This acts as a
- shorthand for ``-fwin-cfg-mechanism=dispatch``.
-- New ``-cl`` option ``/d2guardcfgdispatch-`` added to match MSVC. This acts as a
- shorthand for ``-fwin-cfg-mechanism=check``.
-
-- New option ``-f[no-]strict-bool`` added to control whether Clang can assume
- that ``bool`` values loaded from memory cannot have a bit pattern other
- than 0 or 1.
-
-- New option ``-fcrash-diagnostics-tar`` added to create an archive of crash
- reproducer files for easier bug filing.
-
-- There are a new pair of flags for riscv32 called ``-mzilsd-word-align`` and
- ``-mzilsd-strict-align`` which control whether Zilsd accesses are allowed to
- be aligned to 4-byte alignment rather than fully unaligned or fully (8-byte)
- aligned.
-
-- New ``-cl`` option ``/pathmap:`` added to match MSVC. This option acts as a
- clang's ``-ffile-prefix-map=value`` and has known differences in behaviour
- with the CL's option that do not affect the functionality: nomalizes the
- macro prefix map pathes -- removes `./` and uses the target's platform-
- specific path separator character when expanding the preprocessor macros --
- ``-ffile-reproducible`` (but not the debug and coverage prefix maps);
- does not require ``/experimental:deterministic`` as by MSVC. It needed for
- removing a hostname from a mangling hash gen, but clang-cl does not use
- a hostname when generates the hashes. Known issues -- does not remap the
- source file pathes within PCH/PCM files.
-
-- New ``-cl`` option ``/experimental:deterministic`` added to match CL's option.
- This enables warning emission on usage of non-deterministic macros __DATE__,
- __TIME__ and __TIMESTAMP__ and provides reproducable COFF's timestamp for
- the output object files.
-
-- New ``-cl`` option ``/d1nodatetime`` added to match CL's option. This option
- undefines the standard macros __DATE__, __TIME__ and __TIMESTAMP__ to allow
- reproducable builds. These macros can be redefined from the command line if
- necessary. ``/d1nodatetime-`` can be used to turn this feature off if
- necessary to override the common build settings.
-
-Deprecated Compiler Flags
--------------------------
-
-Modified Compiler Flags
------------------------
-- The `-mno-outline` and `-moutline` compiler flags are now allowed on RISC-V and X86, which both support the machine outliner.
-- The `-mno-outline` flag will now add the `nooutline` IR attribute, so that
- `-mno-outline` and `-moutline` objects can be mixed correctly during LTO.
-- The `-fzero-call-used-regs` compiler flag is now allowed on RISC-V, only the
- "skip", "used-gpr", "used-gpr-arg", "all-gpr" and "all-gpr-arg" options are
- supported for the moment.
-
-- Slightly changed hash id generation to get the unique linkage symbols names
- by ``-unique-internal-linkage-names`` option. Now it uses a path that
- normalized in favor of the target system (same as the preprocessor does
- for the file macros) and allows the reproducable IDs on any build system.
-- ``-fprofile-update=atomic`` will now promote counter updates out of loops,
- similar to the non-atomic case ([#202487](https://github.com/llvm/llvm-project/pull/202487)).
-
-- The ``-cl`` ``/Brepro`` option was modified to match the original CL's option
- and now defines the standard macros __DATE__, __TIME__ and __TIMESTAMP__ to
- "1". The previous functionality remains unchanged.
-
-Removed Compiler Flags
-----------------------
-
-Attribute Changes in Clang
---------------------------
-
-- Added new attribute ``stack_protector_ignore`` to opt specific local variables out of
- the analysis which determines if a function should get a stack protector. A function
- will still generate a stack protector if other local variables or command line flags
- require it.
-
-- Added a new attribute, ``[[clang::no_outline]]`` to suppress outlining from
- annotated functions. This uses the LLVM `nooutline` attribute.
-
-- Introduced a new type attribute ``__attribute__((overflow_behavior))`` which
- currently accepts either ``wrap`` or ``trap`` as an argument, enabling
- type-level control over overflow behavior. There is also an accompanying type
- specifier for each behavior kind via `__ob_wrap` and `__ob_trap`.
-
-- Introduced a new function attribute ``__attribute__((__personality__(...)))``
- to explicitly specify the personality routine for exception handling. THis is
- meant to be a low level tool for language runtime authors to associate a
- foreign language personality with a given function. Note that this does not
- perform any ABI validation for the personality routine.
-
-- :doc:`ThreadSafetyAnalysis` attributes now correctly handle implicit member
- accesses in C, and parameter attributes in C++. This improves diagnostic
- precision and fixes false positives.
-
-- The :doc:`ThreadSafetyAnalysis` attributes ``guarded_by`` and
- ``pt_guarded_by`` now accept multiple capability arguments with refined
- access semantics: *writing* requires all listed capabilities to be held
- exclusively, while *reading* requires at least one to be held. This is
- sound because any writer must hold all capabilities, so holding any one
- prevents concurrent writes.
-
-- :doc:`ThreadSafetyAnalysis` attributes like ``acquire_capability``,
- ``release_capability``, ``requires_capability``, ``locks_excluded``,
- ``try_acquire_capability``, and ``assert_capability`` can now be applied to
- function pointer variables and fields. The analysis checks calls through
- annotated function pointers the same way it checks direct function calls.
- Only plain function pointers are supported; pointers-to-member functions,
- blocks, or wrappers (e.g. ``std::function``) are not yet supported.
-
-- The ``[[clang::unsafe_buffer_usage]]`` attribute is now supported in API
- notes. For example:
-
- .. code-block:: yaml
-
- Functions:
- - Name: myUnsafeFunction
- UnsafeBufferUsage: true
-
-- When using ``-Wunsafe-buffer-usage`` without
- ``-fsafe-buffer-usage-suggestions``, warnings are now emitted only
- once per source file. Pre-compiled code (such as PCH or module
- headers) is no longer repeatedly analyzed, as it is analyzed during
- its initial compilation. (Traditionally included headers are still
- analyzed within each translation unit that includes them). This
- behavior matches most of other ``-W`` diagnostics.
-
- When ``-fsafe-buffer-usage-suggestions`` is enabled, the behavior
- remains the same as before: pre-compiled code is deserialized and
- analyzed alongside the translation unit that uses it, because fix-it
- suggestion analysis requires full visibility of the translation
- unit.
-
-- Added support for ``[[msvc::forceinline]]`` for functions and
- ``[[msvc::forceinline_calls]]`` for statements. Both are aliases to
- ``[[clang::always_inline]]`` with additional checks to ensure that they
- are only accepted in places where MSVC also does.
-
-- The AMDGPU ``amdgpu_num_sgpr`` and ``amdgpu_num_vgpr`` attributes are now
- deprecated. Clang emits a ``-Wdeprecated-declarations`` warning when they
- are used. Use ``amdgpu_waves_per_eu`` instead to control SGPR and VGPR
- usage.
-
-- Clang now allows GNU attributes between a member declarator and bit-field width. (#GH184954)
-
-- The ``[[clang::noescape]]`` attribute now disallows deallocating memory
- through the annotated parameter. This information is currently not exposed to
- LLVM for optimization purposes, to prevent breaking existing adopters. It may
- instead be used by warnings and static analyses to provide more information
- about pointer lifetimes. It may be used to power optimizations in the future,
- however there are no concrete plans to do so at the moment.
-
-- The attributes ``[[clang::opencl_global_device]]`` and ``[[clang::opencl_global_host]]``
- are now deprecated. Clang emits a ``-Wdeprecated-attributes`` warning when
- they are used.
-
-- The ``modular_format`` attribute now supports the ``fixed`` aspect for C
- ISO 18037 fixed-point ``printf`` specifiers.
-
-- The ``const`` and ``pure`` attributes only apply to functions; they are now
- diagnosed and ignored when applied to anything else. Additionally, calling
- a function marked ``noreturn`` from a function marked ``const`` or ``pure``
- is now diagnosed as undefined behavior (#GH129022).
-
-Improvements to Clang's diagnostics
------------------------------------
-- Fixed bug in ``-Wdocumentation`` so that it correctly handles explicit
- function template instantiations (#64087).
-
-- Fixed concept template parameters not being recognized in ``-Wdocumentation``
- when mentioned in tparam comments. (#GH64087)
-
-- ``-Wunused-but-set-variable`` now diagnoses file-scope variables with
- internal linkage (``static`` storage class) that are assigned but never used.
- This new coverage is added under the subgroup ``-Wunused-but-set-global``,
- allowing it to be disabled independently with ``-Wno-unused-but-set-global``.
- (#GH148361)
-
-- Added ``-Wlifetime-safety`` to enable lifetime safety analysis,
- a CFG-based intra-procedural analysis that detects use-after-free and related
- temporal safety bugs. See the
- `RFC <https://discourse.llvm.org/t/rfc-intra-procedural-lifetime-analysis-in-clang/86291>`_
- for more details. By design, this warning is enabled in ``-Weverything``. To disable
- the analysis, use ``-Wno-lifetime-safety`` or ``-fno-lifetime-safety``.
-
-- Added ``-Wlifetime-safety-suggestions`` to enable lifetime annotation suggestions.
- This provides suggestions for function parameters that
- should be marked ``[[clang::lifetimebound]]`` based on lifetime analysis. For
- example, for the following function:
-
- .. code-block:: c++
-
- int* p(int *in) { return in; }
-
- Clang will suggest:
-
- .. code-block:: c++
-
- warning: parameter in intra-TU function should be marked [[clang::lifetimebound]]
- int* p(int *in) { return in; }
- ^~~~~~~
- [[clang::lifetimebound]]
- note: param returned here
- int* p(int *in) { return in; }
- ^~
-
-- Added ``-Wlifetime-safety-noescape`` to detect misuse of ``[[clang::noescape]]``
- annotation where the parameter escapes through return. For example:
-
- .. code-block:: c++
-
- int* p(int *in [[clang::noescape]]) { return in; }
-
- Clang will warn:
-
- .. code-block:: c++
-
- warning: parameter is marked [[clang::noescape]] but escapes
- int* p(int *in [[clang::noescape]]) { return in; }
- ^~~~~~~
- note: returned here
- int* p(int *in [[clang::noescape]]) { return in; }
- ^~
-
-- Added ``-Wlifetime-safety-dangling-field`` to detect dangling field references
- when stack memory escapes to class fields. This is part of ``-Wlifetime-safety``
- and detects cases where local variables or parameters are stored in fields but
- outlive their scope. For example:
-
- .. code-block:: c++
-
- struct DanglingView {
- std::string_view view;
- DanglingView(std::string s) : view(s) {} // warning: address of stack memory escapes to a field
- };
-
-- Improved ``-Wassign-enum`` performance by caching enum enumerator values. (#GH176454)
-
-- Fixed a false negative in ``-Warray-bounds`` where the warning was suppressed
- when accessing a member function on a past-the-end array element.
- (#GH179128)
-
-- Added a missing space to the FixIt for the ``implicit-int`` group of diagnostics and
- made sure that only one such diagnostic and FixIt is emitted per declaration group. (#GH179354)
-
-- Fixed the Fix-It insertion point for ``expected ';' after alias declaration``
- when parsing alias declarations involving a token-split ``>>`` sequence
- (for example, ``using A = X<int>>;``). (#GH184425)
-
-- Fixed incorrect ``implicitly deleted`` diagnostic for explicitly deleted
- candidate function. (#GH185693)
-
-- The ``-Wloop-analysis`` warning has been extended to catch more cases of
- variable modification inside lambda expressions (#GH132038).
-
-- Clang now emits ``-Wsizeof-pointer-memaccess`` when snprintf/vsnprintf use the sizeof
- the destination buffer(dynamically allocated) in the len parameter(#GH162366)
-
-- Added ``-Wmodule-map-path-outside-directory`` (off by default) to warn on
- header and umbrella directory paths that use ``..`` to refer outside the module
- directory in module maps found via implicit search
- (``-fimplicit-module-maps``). This does not affect module maps specified
- explicitly via ``-fmodule-map-file=``.
-
-- Honour ``[[maybe_unused]]`` attribute on private fields.
- ``-Wunused-private-field`` no longer emits a warning for annotated private
- fields.
-
-- Improved ``-Wgnu-zero-variadic-macro-arguments`` to suggest using
- ``__VA_OPT__`` if the current language version supports it(#GH188624)
-
-- Clang now emits an error when implicitly casting a complex type to a built-in vector type. (#GH186805)
-
-- Added ``-Wnonportable-include-path-separator`` (off by default) to catch
- #include directives that use backslashes as a path separator. The warning
- includes a FixIt to change all the backslashes to forward slashes, so that the
- code can automatically be made portable to other host platforms that don't
- support backslashes.
-
-- Clang now explains why template deduction fails for explicit template arguments.
-
-- No longer emitting a ``-Wpre-c2y-compat`` or extension diagnostic about use
- of octal literals with a ``0o`` prefix, and no longer emitting a
- ``-Wdeprecated-octal-literals`` diagnostic for use of octal literals without
- a ``0o`` prefix, when the literal is expanded from a macro defined in a
- system header. (#GH192389)
-
-- Improved error recovery for missing semicolons after class members. Clang now avoids
- skipping subsequent valid declarations when their previous decl is missing semicolon.
-
-- Removed the body of lambdas from some diagnostic messages.
-
-- Fixed false positive host-device mismatch errors in discarded `if constexpr` branches for CUDA/HIP;
- such calls are now correctly skipped.
-
-- Clang now errors when a function declaration aliases a variable or vice versa. (#GH195550)
-
-- Added ``-Wattribute-alias`` to diagnose type mismatches between an alias and its aliased function. (#GH195550)
-
-- The diagnostics around ``__block`` now explain why a variable cannot be marked ``__block``. (#GH197213)
-
-- Extended ``-Wnonportable-include-path`` to warn about trailing whitespace and dots in ``#include`` paths. (#GH190610)
-
-- Clang now emits error when attribute is missing closing ``]]`` followed by ``;;``. (#GH187223)
-
-- Clang now rejects inline asm constraints and clobbers that contain an
- embedded null character, instead of silently truncating them. (#GH173900)
-
-- Added ``-Wstringop-overread`` to warn when ``memcpy``, ``memmove``, ``memcmp``,
- and related builtins read more bytes than the source buffer size (#GH83728).
-
-- Diagnostics for the C++11 range-based for statement now report the correct
- iterator type in notes for invalid iterator types.
-
-Improvements to Clang's time-trace
-----------------------------------
-
-Improvements to Coverage Mapping
---------------------------------
-
-- [MC/DC] Nested expressions are handled as individual MC/DC expressions.
-- "Single byte coverage" now supports branch coverage and can be used
- together with ``-fcoverage-mcdc``.
-- Consteval member functions are no longer emitted in coverage mappings,
- matching the existing behavior for free consteval functions. (#GH164448)
-
-Bug Fixes in This Version
--------------------------
-
-- Fixed an assertion when comparing a fixed point type with a ``_BitInt`` type. (#GH196948)
-- Fixed atomic boolean compound assignment; the conversion back to atomic bool would be miscompiled. (#GH33210)
-- Correctly handle default template argument when establishing subsumption. (#GH188640)
-- Fixed a failed assertion in the preprocessor when ``__has_embed`` parameters are missing parentheses. (#GH175088)
-- Fix lifetime extension of temporaries in for-range-initializers in templates. (#GH165182)
-- Fixed a preprocessor crash in ``__has_cpp_attribute`` on incomplete scoped attributes. (#GH178098)
-- Fixes an assertion failure when evaluating ``__underlying_type`` on enum redeclarations. (#GH177943)
-- Fixed an assertion failure caused by nested macro expansion during header-name lexing (``__has_embed(__has_include)``). (#GH178635)
-- Clang now outputs relative paths of embeds for dependency output. (#GH161950)
-- Fix the result type of a binary operation where both operands are 'void' l-values. (#GH111300)
-- Fixed an assertion failure when evaluating ``_Countof`` on invalid ``void``-typed operands. (#GH180893)
-- Fixed an assertion failure in the serialized diagnostic printer when it is destroyed without calling ``finish()``. (#GH140433)
-- Fixed an assertion failure caused by error recovery while extending a nested name specifier with results from ordinary lookup. (#GH181470)
-- Fixed a crash when parsing ``#pragma clang attribute`` arguments for attributes that forbid arguments. (#GH182122)
-- Fixed a bug in how Clang re-transforms expressions produced from substititions
- from type aliases and concept specializations. (#GH191738) (#GH196375)
-- Fixed a bug with multiple-include optimization (MIOpt) state not being preserved in some cases during lexing, which could suppress header-guard mismatch diagnostics and interfere with include-guard optimization. (#GH180155)
-- Fixed a crash when normalizing constraints involving concept template parameters whose index coincided with non-concept template parameters in the same parameter mapping.
-- Fixed a crash caused by accessing dependent diagnostics of a non-dependent context.
-- Fixed a crash when substituting into a non-type template parameter that has a type containing an undeduced placeholder type.
-- Fixed several crashes and improved diagnostics when a multidimensional subscript operator is applied to a built-in type. (#GH187800)
-- Correctly diagnosing and no longer crashing when ``export module foo``
- (without a semicolon) are the final tokens in a module file. (#GH187771)
-- Fixed a crash in duplicate attribute checking caused by comparing constant arguments with different integer signedness. (#GH188259)
-- Clang now emits an error when returning an initializer list from a lambda
- with an explicit return type of void. The diagnostic now correctly refers
- to "lambda" instead of "block". (#GH188661)
-- Fixed a crash on _BitInt(N) arrays where 129 ≤ N ≤ 192 due to incorrect array filler lowering. (#GH189643)
-- Fixed the behavior in C23 of ``auto``, by emitting an error when an array type is specified for a ``char *``. (#GH162694)
-- Fixed an issue where an assert was thrown instead of an error if no vulkan env was specified with ``--triple spirv``. (#GH189964)
-- Fixed incorrect rejection of ``auto`` with reordered declaration specifiers in C23. (#GH164121)
-- Fixed a bug where Clang fails to find instantiation of Decls in constraint checking. (#GH173086)
-- Fixed a crash when assigning to an element of an ``ext_vector_type`` with ``bool`` element type. (#GH189260)
-- Fixed a crash caused by declaring multiple ``ownership_returns`` attributes with mismatched or missing arguments. (#GH188733)
-- Clang now emits an error for friend declarations of lambda members. (#GH26540)
-- Fixed a crash caused by lambda capture handling in delayed default arguments. (#GH176534)
-- Fixed a crash when parsing invalid ``static_assert`` declarations with string-literal messages (#GH187690).
-- Fixed a potential stack-use-after-return issue in Clang when copy-initializing
- an array via an element-at-a-time copy loop (#GH192026)
-- Fixed an issue where certain designated initializers would be rejected for constexpr variables. (#GH193373)
-- Fixed ``clang::Preprocessor::recomputeCurLexerKind`` to avoid default fallback to ``CurLexerCallback = CLK_CachingLexer;``. This prevents code-completion
- EOF handling from accidentally restoring CLK_CachingLexer while a tentative parse is still active, which could trigger a caching lexer re-entry assertion
- in clangd signature help. (#GH200677)
-- Fixed a crash when ``#embed`` is used with C++ modules (#GH195350)
-- Fixed an assertion in constant evaluation when using a defaulted comparison operator in a ``union``. (#GH147127)
-- Fixed a bug where ``-x cuda`` caused clang to immediately resolve templates that should not be. (#GH200545)
-- Fixed an issue where ``__typeof_unqual`` and ``__typeof_unqual__`` were rejected as a declaration specifier in block scope in C++.
-- Fixed crash when checking for overflow for unary operator that can't overflow (#GH170072)
-- Clang no longer handles a `" q-char-sequence "` header name as a string literal (#GH132643).
-- Fixed an assertion where we improperly handled implicit conversions to integral types from an atomic-type with a conversion function. (#GH201770)
-- Fixed assertion failures involving code completion with delayed default arguments and exception specifications. (#GH200879)
-- Fixed a regression where calling a function that takes a class-type parameter by value inside ``decltype`` of a concept could be incorrectly rejected when used as a non-type template argument. (#GH175831)
-
-Bug Fixes to Compiler Builtins
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-- Fix a crash when passing an unresolved overload set to ``__builtin_classify_type``. (#GH175589)
-- Fixed a crash when calling `__builtin_allow_sanitize_check` with no arguments. (#GH183927)
-- ``__annotation`` is now diagnosed as unsupported on non-Windows/UEFI targets, fixing a
- crash when using it with ``-fms-extensions`` on other platforms. (#GH184318)
-- Fixed a compiler crash due to an unresolved overloaded function type when
- calling ``__builtin_bit_cast``. (#GH200112)
-
-Bug Fixes to Attribute Support
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-- Fixed a behavioral discrepancy between deleted functions and private members when checking the ``enable_if`` attribute. (#GH175895)
-- Fixed ``init_priority`` attribute by delaying type checks until after the type is deduced.
-- Fixed a crash when a ``section`` attribute or ``#pragma clang section`` caused a
- section type conflict with a declaration whose name is not a simple identifier,
- such as a lambda's call operator. (#GH192264)
-- Fixed a regression where attributed types (such as those carrying ``_Nonnull``/``_Nullable`` attributes)
- were not deduplicated, because the attributes' arguments were not taken into
- account when uniquing them. The duplications could substantially increase the
- size of precompiled headers and modules (PCH/PCM), and the time spent loading
- them. (#GH200961)
-
-Bug Fixes to C++ Support
-^^^^^^^^^^^^^^^^^^^^^^^^
-- Fixed a preprocessor assertion failure triggered when parsing an invalid template-id starting with ``::template operator``. (#GH186582)
-- Fixed a crash when a function template is defined as a non-template friend with a global scope qualifier. (#GH185341)
-- Clang now rejects constant template parameters with block pointer types, since these are not implemented anyway and would lead to crashes. (#GH189247)
-- Fixed a crash on error recovery when dealing with invalid templates. (#GH183075)
-- Fixed a crash when instantiating ``requires`` expressions involving substitution failures in C++ concepts. (#GH176402)
-- Concepts appearing in the require-clause of a member function no longer have access to non-public members of that class,
- or to a current class object. (#GH115838) (#GH194803) (#GH197067)
-- We no longer caches invalid variable specializations. (#GH132592)
-- Fixed an incorrect template argument deduction when matching packs of template
- template parameters when one of its parameters is also a pack. (#GH181166)
-- Clang no longer errors on overloads with different ref-qualifiers and constraints. (#GH120812)
-- Fixed a crash when a default argument is passed to an explicit object parameter. (#GH176639)
-- Fixed an alias template CTAD crash.
-- Correctly diagnose uses of ``co_await`` / ``co_yield`` in the default argument of nested function declarations. (#GH98923)
-- Fixed a crash when diagnosing an invalid static member function with an explicit object parameter (#GH177741)
-- Clang incorrectly instantiated variable specializations outside of the immediate context. (#GH54439)
-- Fixed a crash when pack expansions are used as arguments for non-pack parameters of built-in templates. (#GH180307)
-- Fixed crash instantiating class member specializations.
-- Fix a problem where a substitution failure when evaluating a type requirement
- could directly make the program ill-formed.
-- Fix a problem where pack index expressions where incorrectly being regarded as equivalent.
-- Fixed a bug where captured variables in non-mutable lambdas were incorrectly treated as mutable
- when used inside decltype in the return type. (#GH180460)
-- Fixed a crash when evaluating uninitialized GCC vector/ext_vector_type vectors in ``constexpr``. (#GH180044)
-- Fixed a crash when `explicit(bool)` is used with an incomplete enumeration. (#GH183887)
-- Fixed a crash on ``typeid`` of incomplete local types during template instantiation. (#GH63242), (#GH176397)
-- Fixed spurious diagnostics produced when checking if constraints are equivalent for redeclarations,
- which could make the program mistakenly ill-formed.
-- Fixed a crash when an immediate-invoked ``consteval`` lambda is used as an invalid initializer. (#GH185270)
-- Fixed an assertion failure when using a global destructor with a target with a non-default program address space. (#GH186484)
-- Fixed a crash when instantiating an invalid out-of-line static data member definition in a local class. (#GH176152)
-- Inherited constructors in ``dllexport`` classes are now exported for ABI-compatible cases, matching
- MSVC behavior. Constructors with variadic arguments or callee-cleanup parameters are not yet supported
- and produce a warning. (#GH162640)
-- Correctly diagnose invalid non-dependent calls in dependent contexts. (#GH135694)
-- Fix initialization of GRO when GRO-return type mismatches, as part of CWG2563. (#GH98744)
-- Fix an error using an initializer list with array new for a type that is not default-constructible. (#GH81157)
-- We no longer consider conversion operators when copy-initializing from the same type. This was non
- conforming and could lead to recursive constraint satisfaction checking. (#GH149443)
-- Fixed a crash in Itanium C++ name mangling for a lambda in a local class field initializer inside a constructor/destructor. (#GH176395)
-- Fixed a crash when Expr::ClassifyImpl computes a classification like CL_LValue or CL_PRValue, then asserts that this
- agrees with the AST node's own value category. (#GH202693)
-- Fixed crashes in Itanium C++ name mangling for lambdas with trailing requires-clauses involving requires-expressions. (#GH100774) (#GH123854)
-- Fixed an invalid rejection and assertion failure while generating ``operator=`` for fields with the ``__restrict`` qualifier. (#GH37979)
-- Fixed a use-after-free bug when parsing default arguments containing lambdas in declarations with template-id declarators. (#GH196725)
-- Fixed a crash in constant evaluation using placement new on an array which was later initialized. (#GH196450)
-- Fixed an issue where Clang incorrectly accepted invalid unqualified uses of local nested class names outside their declaring scope. (#GH184622)
-- Fixed a crash when parsing invalid friend declaration with storage-class specifier. (#GH186569)
-- Fixed a missing vtable for ``dynamic_cast<FinalClass *>(this)`` in a function template. (#GH198511)
-- Fixed an assertion failure during init-list checking of an array whose element type is an incomplete class. (#GH140685)
-
-Bug Fixes to AST Handling
-^^^^^^^^^^^^^^^^^^^^^^^^^
-- Fixed a bug where explicit nullability property attributes were not stored in AST nodes in Objective-C. (#GH179703)
-- Fixed a bug where alias CTAD, or an invalid template template parameter, could create a template with an empty template
- parameter list. This also adds asserts to prevent this from happening again.
-- Fixed a crash when parsing Doxygen ``@param`` commands attached to invalid declarations or non-function entities. (#GH182737)
-- Fixed the SourceLocation and SourceRange of reversed rewritten CXXOperatorCallExpr. (#GH192467)
-- Fixed a assertion when ``__block`` is used on global variables in C mode. (#GH183974)
-- Added missing AST nodes representing the ``decltype`` specifiers in destructor call to AST.
-- Fixed a missing ODR violation diagnostic introduced by the inline assembly string or clobber list. (#GH198616)
-
-Miscellaneous Bug Fixes
-^^^^^^^^^^^^^^^^^^^^^^^
-- Fixed a crash whith the AST text dumper, when dumping a reference to a
- decomposition with no bindinds. (#GH198842)
-- Fixed the arguments of the format attribute on ``__builtin_os_log_format``. Previously, they were off by 1.
-
-Miscellaneous Clang Crashes Fixed
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-
-- Fixed a crash when attempting to jump over initialization of a variable with variably modified type. (#GH175540)
-- Fixed a crash when using loop hint with a value dependent argument inside a
- generic lambda. (#GH172289)
-- Fixed a crash in C++ overload resolution with ``_Atomic``-qualified argument types. (#GH170433)
-- Fixed a crash related to missing source locations (#GH186655)
-- Fixed a crash when casting a parenthesized unresolved template-id or array section. (#GH183505)
-- Fixed a crash when initializing a ``constexpr`` pointer with a floating-point literal in C23. (#GH180313)
-- Fixed a lack of diagnostic for substitution failures in base classes when using `std::void_t`-like types.
-- Fixed a crash when emitting debug info for base classes with instantiation-dependent-only types (#GH193932)
-- Fixed an assertion when diagnosing address-space qualified ``new``/``delete`` in language-defined address spaces such as OpenCL ``__local``. (#GH178319)
-- Fixed an assertion failure in ObjC++ ARC when binding a rvalue reference to reference with different lifetimes (#GH178524)
-- Fixed a crash when subscripting a vector type with large unsigned integer values. (#GH180563)
-- Fixed a crash when attempting to diagnose incompatible conversions involving function types (#GH182534)
-- Fixed a crash when evaluating ``__is_bitwise_cloneable`` on invalid record types. (#GH183707)
-- Fixed an assertion failure when casting a function pointer with a target with a non-default program address space. (#GH186210)
-- Fixed a crash when ``decltype(__builtin_FUNCTION())`` is used as a template type argument. (#GH167433)
-- Fixed an assertion failure when parsing an invalid ``decltype`` specifier with missing parentheses or extra semicolons. (#GH188014)
-- Fixed a crash when explicitly casting a complex type to or from an atomic complex type. (#GH172208)
-- Fixed a crash when explicitly casting a scalar to an atomic complex. (#GH114885)
-- Fixed an assertion failure when parsing an invalid out-of-line enum definition with template parameters. (#GH187909)
-- Fixed an assertion failure on invalid template template parameter during typo correction. (#GH183983)
-- Fixed an assertion failure in ``isAtEndOfMacroExpansion`` on macro expansions crossing the boundary of two fileIDs. (#GH115007), (#GH21755)
-- Fixed an assertion failure when ``__builtin_dump_struct`` is used with an
- immediate-escalated callable. (#GH192846)
-- Fixed a crash when passing one sized implicitly casted vector to a ``abs`` function. (#GH204777)
-- Fixed a crash when diagnosing an invalid out-of-line definition of a member class template. (#GH201490)
-
-OpenACC Specific Changes
-------------------------
-
-OpenCL Specific Changes
------------------------
-- Added support for OpenCL C 3.1 language version (``-cl-std=CL3.1``).
-
-Target Specific Changes
------------------------
-
-AMDGPU Support
-^^^^^^^^^^^^^^
-- Introduced a new target specific builtin ``__builtin_amdgcn_processor_is``,
- a late / deferred query for the current target processor.
-- Introduced a new target specific builtin ``__builtin_amdgcn_is_invocable``,
- a late / deferred query for the availability of target specific builtins.
-- Initial support for gfx1310
-- The ``amdgpu_num_sgpr`` and ``amdgpu_num_vgpr`` function attributes are now
- deprecated. Using them produces a ``-Wdeprecated-declarations`` warning. Use
- ``amdgpu_waves_per_eu`` instead.
-
-NVPTX Support
-^^^^^^^^^^^^^^
-
-X86 Support
-^^^^^^^^^^^
-- ``march=znver6`` is now supported.
-
-Arm and AArch64 Support
-^^^^^^^^^^^^^^^^^^^^^^^
-
-- Support has been added for the following processors (-mcpu identifiers in parenthesis):
-
-* Arm AGI CPU (armagicpu).
-* Hisilicon hip12 core (hip12).
-
-Android Support
-^^^^^^^^^^^^^^^
-
-Windows Support
-^^^^^^^^^^^^^^^
-
-- Clang now defines the ``_MSVC_TRADITIONAL`` macro as ``1`` when emulating MSVC
- 19.15 (Visual Studio 2017 version 15.8) and later. (#GH47114)
-- ``-fmacro-prefix-map=`` (``-ffile-prefix-map=``) now affects an anonymous namespace hash generation
- for the MSVC targets and allows deterministic symbol mangling for reproducible builds.
-
-- Added the ``-fwinx64-eh-unwind=`` flag to select the x64 Windows unwind info
- version (``v1``, ``v2-best-effort``, ``v2-required``, or ``v3``). The legacy
- ``-fwinx64-eh-unwindv2=`` flag is deprecated; it is still accepted and mapped
- onto the new flag as follows:
-
- .. list-table::
- :header-rows: 1
-
- * - Legacy ``-fwinx64-eh-unwindv2=``
- - New ``-fwinx64-eh-unwind=``
- * - ``disabled``
- - ``v1`` (default; no flag forwarded)
- * - ``best-effort``
- - ``v2-best-effort``
- * - ``required``
- - ``v2-required``
-
- The MSVC-compatible ``/d2epilogunwind`` and ``/d2epilogunwindrequirev2``
- options map to ``v2-best-effort`` and ``v2-required`` respectively.
-
-- When targeting Windows x64 with EGPR (`-mapx-features=egpr`), Clang now
- automatically enables V3 unwind info (`-fwinx64-eh-unwind=v3`) if no
- explicit unwind version was specified.
-
-LoongArch Support
-^^^^^^^^^^^^^^^^^
-
-- DWARF fission is now compatible with linker relaxations, allowing `-gsplit-dwarf` and `-mrelax`
- to be used together when building for the LoongArch platform.
-
-RISC-V Support
-^^^^^^^^^^^^^^
-
-- Tenstorrent Ascalon D8 was renamed to Ascalon X. Use `tt-ascalon-x` with `-mcpu` or `-mtune`.
-- Intrinsics were added for the 'Zvabd` (RISC-V Integer Vector Absolute Difference) extension.
-- Intrinsics were added for the 'Zvzip` (Reordering Structured Data in Vector Registers) extension.
-- A new ``-mtune`` syntax was added to support processor-specific tuning feature string
- Currently this new syntax is gated by the ``-mexperimental-mtune-syntax`` flag.
-
-
-CUDA/HIP Language Changes
-^^^^^^^^^^^^^^^^^^^^^^^^^
-
-- The new offloading driver is now the default for HIP. Use
- `--no-oflfoad-new-driver` to return to the old behavior.
-
-CUDA Support
-^^^^^^^^^^^^
-
-- Fixed a bug where host-device ambiguities in CUDA/HIP when retrieving the
- address of specializations of templated functions that have overloads for both
- host and device. (#GH199299)
-
-AIX Support
-^^^^^^^^^^^
-
-- The driver default for the linker flag `-bcdtors` now defaults to `mbr`
- (instead of `all`) which only extracts static init from archive members which
- would otherwise be referenced.
- (See https://www.ibm.com/docs/en/aix/7.2.0?topic=l-ld-command for details).
-- The driver now uses ``-lcompiler_rt`` instead of ``-latomic``, and the compiler-rt
- archive has been renamed from ``libatomic.a`` to ``libcompiler_rt.a`` to avoid conflicts
- between the LLVM libatomic and the GNU libatomic from the AIX toolbox as they share
- the same library name.
-- Added support for ``#pragma comment(copyright, "token_sequence")`` on AIX.
- This directive embeds a copyright or identifying string into the compiled object file.
- The string is included in the final executable and loaded into memory at program runtime.
-- The driver relaxes the restrictions on the ``OBJECT_MODE`` environment
- variable and now silently accepts ``32_64`` and ``any``.
-
-NetBSD Support
-^^^^^^^^^^^^^^
-
-WebAssembly Support
-^^^^^^^^^^^^^^^^^^^
-
-- Fixed a crash when ``__funcref`` is applied to a non-function pointer type.
- (#GH118233)
-- WebAssembly reference types (``__externref_t`` and ``__funcref`` function
- pointers) now lower to the opaque IR types ``target("wasm.externref")`` and
- ``target("wasm.funcref")`` instead of ``ptr addrspace(10)`` /
- ``ptr addrspace(20)``.
-- Fixed a compiler crash at ``-O2`` when reference-type values were passed
- through control flow that the SLP vectorizer tried to vectorize.
-
-AVR Support
-^^^^^^^^^^^
-
-SystemZ Support
-^^^^^^^^^^^^^^^
-
-- Add support for `#pragma export` for z/OS. This is a pragma used to export functions and variables
- with external linkage from shared libraries. It provides compatibility with the IBM XL C/C++
- compiler.
-
-DWARF Support in Clang
-----------------------
-
-Floating Point Support in Clang
--------------------------------
-
-Fixed Point Support in Clang
-----------------------------
-
-AST Matchers
-------------
-- Add ``functionTypeLoc`` matcher for matching ``FunctionTypeLoc``.
-- Add missing support for ``TraversalKind`` in some ``addMatcher()`` overloads.
-
-clang-format
-------------
-- Add ``ObjCSpaceAfterMethodDeclarationPrefix`` option to control space between the
- '-'/'+' and the return type in Objective-C method declarations
-- Deprecate the ``BinPackParameters`` and ``BinPackArguments`` options and replace
- them with the ``PackParameters`` and ``PackArguments`` structs (respectively) to
- unify packing behavior. Add the ``BreakAfter`` option to the structs, allowing
- parameter and argument lists to be formatted with one parameter/argument on each
- line if they exceed the specified count.
-- Add ``AfterComma`` value to ``BreakConstructorInitializers`` to allow breaking
- constructor initializers after commas, keeping the colon on the same line.
-- Extend ``BreakBinaryOperations`` to accept a structured configuration with
- per-operator break rules and minimum chain length gating via ``PerOperator``.
-- Add ``AllowShortRecordOnASingleLine`` option and set it to ``EmptyAndAttached`` for LLVM style.
-- Add ``BreakFunctionDeclarationParameters`` option to always break before function
- declaration parameters.
-- Add ``EnumAssignments`` option to ``AlignConsecutiveAssignments`` for aligning
- enum assignments without affecting other assignments.
-- Add ``BreakBeforeReturnType`` option to break before the function return
- type.
-
-libclang
---------
-- Visit constraints of `auto` type to properly visit concept usages (#GH166580)
-- Visit switch initializer statements (https://bugs.kde.org/show_bug.cgi?id=415537#c2)
-- Fix crash in clang_getBinaryOperatorKindSpelling and clang_getUnaryOperatorKindSpelling
-- The clang_Module_getASTFile API is deprecated and now always returns nullptr
-- The clang_Cursor_getCommentRange API will now return a comment range for macro definitions that have documentation comments.
-- Added CXType_PredefinedSugar for __ptrdiff_t, __size_t, and
- __signed_size_t types, which are no longer exposed as
- CXType_Unexposed.
-
-Code Completion
----------------
-
-- Fixed a crash in code completion when using a C-Style cast with a parenthesized
- operand in Objective-C++ mode. (#GH180125)
-
-Static Analyzer
----------------
-
-Crash and bug fixes
-^^^^^^^^^^^^^^^^^^^
-
-- Fixed ``security.VAList`` checker producing false positives when analyzing
- C23 code where ``va_start`` expands to ``__builtin_c23_va_start``.
-
-- Fixed a compiler crash when combining ``_Atomic`` and ``__auto_type``
- in C, for example ``_Atomic __auto_type x = expr``. (#GH118058)
-
-Improvements
-^^^^^^^^^^^^
-
-- ``alpha.unix.PthreadLock`` now emits path notes on lock, unlock, destroy,
- and init operations.
-
-.. comment:
- This is for the Static Analyzer.
- Using the caret `^^^` underlining for subsections:
- - Crash and bug fixes
- - New checkers and features
- - Improvements
- - Moved checkers
-
-
-Moved checkers
-^^^^^^^^^^^^^^
-
-- The checker ``unix.cstring.UninitializedRead`` is now out of alpha.
-
-.. _release-notes-sanitizers:
-
-Sanitizers
-----------
-- UndefinedBehaviorSanitizer now supports ``__ubsan_default_suppressions``.
-
-- Sanitizer Special Case Lists (``-fsanitize-ignorelist``) now support
- Version 4 of the Special Case List format, which introduces a transition
- period for leading dot-slash (``./``) canonicalization in path matching.
- Version 4 matches both canonicalized and non-canonicalized paths but emits a
- warning for deprecated matches. Version 5 drops backward compatibility and
- requires rules to match canonicalized paths (without leading ``./``).
-
-- Sanitizer Special Case Lists (``-fsanitize-ignorelist``) and warning
- suppression mappings (``--warning-suppression-mappings``) now recognize version
- 4 of the Special Case List format (indicated by ``#!special-case-list-v4``).
- On Windows hosts, path matching is slash-agnostic (both forward slashes (``/``)
- and backslashes (``\``) match either path separator in both patterns and paths).
-
-Python Binding Changes
-----------------------
-- Add deprecation warnings to ``CompletionChunk.isKind...`` methods.
- These will be removed in a future release. Existing uses should be adapted
- to directly compare equality of the ``CompletionChunk`` kind with
- the corresponding ``CompletionChunkKind`` variant.
-
- Affected methods: ``isKindOptional``, ``isKindTypedText``, ``isKindPlaceHolder``,
- ``isKindInformative`` and ``isKindResultType``.
-- Add a deprecation warning to ``CodeCompletionResults.results``.
- This property will become an implementation detail with changed behavior in a
- future release and should not be used directly.. Existing uses of
- ``CodeCompletionResults.results`` should be changed to directly use
- ``CodeCompletionResults``: it nows supports ``__len__`` and ``__getitem__``,
- so it can be used the same as ``CodeCompletionResults.results``.
-- Added a new helper method ``get_clang_version`` to the class ``Config`` to
- read the version string of the libclang in use.
-
-OpenMP Support
---------------
-- Added support for ``transparent`` clause in task and taskloop directives.
-- Added support for ``use_device_ptr`` clause to accept an optional
- ``fallback`` modifier (``fb_nullify`` or ``fb_preserve``) with OpenMP >= 61.
-- Added support for ``local`` clause with declare_target directive when
- OpenMP >= 60.
-- Fixed the identity element used for ``reduction(* : x)`` over C++ class types
- (e.g. ``std::complex``). The private copy is now initialized to the
- multiplicative identity instead of being value-initialized, which previously
- produced a wrong result (the product collapsed to the additive identity).
-
-SYCL Support
-------------
-- SYCL compilations now default to ``-std=c++17`` when no explicit language
- standard is specified. Standards below C++17 are rejected with a diagnostic.
-
-- Clang now assumes default target for SYCL device compilation is 64-bit SPIR-V
- and it now diagnoses if a non-supporting target is specified via command line.
- (#GH167358)
-
-- The SYCL runtime shared library has been renamed from ``libsycl.so`` to
- ``libLLVMSYCL.so`` to align with LLVM naming conventions.
-
-- SYCL header include paths are now added automatically for both host and
- device compilations.
-
-- SYCL runtime library linking is now supported on Windows. When ``-fsycl`` is
- specified, Clang automatically adds ``/MD`` if no explicit CRT flag is
- present, links the appropriate debug (``LLVMSYCLd.lib``) or release
- (``LLVMSYCL.lib``) library, and rejects static CRT flags (``/MT``,
- ``/MTd``) with a diagnostic. Use ``-nolibsycl`` to suppress automatic
- library linking.
-
-- Fixed ``-nolibsycl`` being silently ignored on Linux: the SYCL runtime
- library was unconditionally added to the link line even when the flag was
- passed.
-
-- Clang now is capable of diagnosing reference kernel parameters which are not
- allowed by SYCL 2020 spec.
-
-Improvements
-^^^^^^^^^^^^
-- Improved substitution performance in concept checking. (#GH172266)
-
-- Clang now preserves the left-hand side of a binary expression (such as an
- assignment or comparison) in a ``RecoveryExpr`` when the right-hand side fails
- to parse. This improves IDE features like go-to-definition in ``clangd``.
-
-Additional Information
-======================
-
-A wide variety of additional information is available on the `Clang web
-page <https://clang.llvm.org/>`_. The web page contains versions of the
-API documentation which are up-to-date with the Git version of
-the source code. You can access versions of these documents specific to
-this release by going into the "``clang/docs/``" directory in the Clang
-tree.
-
-If you have any questions or comments about Clang, please feel free to
-contact us on the `Discourse forums (Clang Frontend category)
-<https://discourse.llvm.org/c/clang/6>`_.
More information about the cfe-commits
mailing list