diff --git a/clang-tools-extra/clang-tidy/utils/RenamerClangTidyCheck.cpp b/clang-tools-extra/clang-tidy/utils/RenamerClangTidyCheck.cpp index 69b7d40ef628d..ad8048e2a92b7 100644 --- a/clang-tools-extra/clang-tidy/utils/RenamerClangTidyCheck.cpp +++ b/clang-tools-extra/clang-tidy/utils/RenamerClangTidyCheck.cpp @@ -489,7 +489,7 @@ void RenamerClangTidyCheck::checkNamedDecl(const NamedDecl *Decl, } Failure.Info = std::move(Info); - addUsage(Decl, Range); + addUsage(Decl, Range, &SourceMgr); } void RenamerClangTidyCheck::check(const MatchFinder::MatchResult &Result) { diff --git a/clang-tools-extra/docs/ReleaseNotes.rst b/clang-tools-extra/docs/ReleaseNotes.rst index a7193e90c38da..b66be44e9f8a6 100644 --- a/clang-tools-extra/docs/ReleaseNotes.rst +++ b/clang-tools-extra/docs/ReleaseNotes.rst @@ -268,7 +268,7 @@ Changes in existing checks ` check in `GetConfigPerFile` mode by resolving symbolic links to header files. Fixed handling of Hungarian Prefix when configured to `LowerCase`. Added support for renaming designated - initializers. + initializers. Added support for renaming macro arguments. - Improved :doc:`readability-implicit-bool-conversion ` check to provide diff --git a/clang-tools-extra/test/clang-tidy/checkers/readability/identifier-naming.cpp b/clang-tools-extra/test/clang-tidy/checkers/readability/identifier-naming.cpp index 57ef4aae5ddb7..99149fe86acee 100644 --- a/clang-tools-extra/test/clang-tidy/checkers/readability/identifier-naming.cpp +++ b/clang-tools-extra/test/clang-tidy/checkers/readability/identifier-naming.cpp @@ -108,10 +108,12 @@ USER_NS::object g_s2; // NO warnings or fixes expected as USER_NS and object are declared in a header file SYSTEM_MACRO(var1); -// NO warnings or fixes expected as var1 is from macro expansion +// CHECK-MESSAGES: :[[@LINE-1]]:14: warning: invalid case style for global variable 'var1' [readability-identifier-naming] +// CHECK-FIXES: {{^}}SYSTEM_MACRO(g_var1); USER_MACRO(var2); -// NO warnings or fixes expected as var2 is declared in a macro expansion +// CHECK-MESSAGES: :[[@LINE-1]]:12: warning: invalid case style for global variable 'var2' [readability-identifier-naming] +// CHECK-FIXES: {{^}}USER_MACRO(g_var2); #define BLA int FOO_bar BLA; @@ -602,9 +604,20 @@ static void static_Function() { // CHECK-FIXES: {{^}}#define MY_TEST_MACRO(X) X() void MY_TEST_Macro(function) {} -// CHECK-FIXES: {{^}}void MY_TEST_MACRO(function) {} -} -} +// CHECK-MESSAGES: :[[@LINE-1]]:20: warning: invalid case style for global function 'function' [readability-identifier-naming] +// CHECK-FIXES: {{^}}void MY_TEST_MACRO(Function) {} + +#define MY_CAT_IMPL(l, r) l ## r +#define MY_CAT(l, r) MY_CAT_IMPL(l, r) +#define MY_MACRO2(foo) int MY_CAT(awesome_, MY_CAT(foo, __COUNTER__)) = 0 +#define MY_MACRO3(foo) int MY_CAT(awesome_, foo) = 0 +MY_MACRO2(myglob); +MY_MACRO3(myglob); +// No suggestions should occur even though the resulting decl of awesome_myglob# +// or awesome_myglob are not entirely within a macro argument. + +} // namespace InlineNamespace +} // namespace FOO_NS template struct a { // CHECK-MESSAGES: :[[@LINE-1]]:32: warning: invalid case style for struct 'a' diff --git a/clang/docs/ReleaseNotes.rst b/clang/docs/ReleaseNotes.rst index f96cebbde3d82..f5359afe1f099 100644 --- a/clang/docs/ReleaseNotes.rst +++ b/clang/docs/ReleaseNotes.rst @@ -520,12 +520,11 @@ Bug Fixes to C++ Support - Fix an issue caused by not handling invalid cases when substituting into the parameter mapping of a constraint. Fixes (#GH86757). - Fixed a bug that prevented member function templates of class templates declared with a deduced return type from being explicitly specialized for a given implicit instantiation of the class template. -- Fixed a crash when ``this`` is used in a dependent class scope function template specialization - that instantiates to a static member function. - Fix crash when inheriting from a cv-qualified type. Fixes: (`#35603 `_) - Fix a crash when the using enum declaration uses an anonymous enumeration. Fixes (#GH86790). +- Handled an edge case in ``getFullyPackExpandedSize`` so that we now avoid a false-positive diagnostic. (#GH84220) - Clang now correctly tracks type dependence of by-value captures in lambdas with an explicit object parameter. Fixes (#GH70604), (#GH79754), (#GH84163), (#GH84425), (#GH86054), (#GH86398), and (#GH86399). diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 790f41627522d..68f7fb48c4dc5 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -185,6 +185,7 @@ class PseudoObjectExpr; class QualType; class SemaHLSL; class SemaOpenACC; +class SemaSYCL; class StandardConversionSequence; class Stmt; class StringLiteral; @@ -295,193 +296,6 @@ class FileNullabilityMap { } }; -// TODO SYCL Integration header approach relies on an assumption that kernel -// lambda objects created by the host compiler and any of the device compilers -// will be identical wrt to field types, order and offsets. Some verification -// mechanism should be developed to enforce that. - -// TODO FIXME SYCL Support for SYCL in FE should be refactored: -// - kernel identification and generation should be made a separate pass over -// AST. RecursiveASTVisitor + VisitFunctionTemplateDecl + -// FunctionTemplateDecl::getSpecializations() mechanism could be used for that. -// - All SYCL stuff on Sema level should be encapsulated into a single Sema -// field -// - Move SYCL stuff into a separate header - -// Represents contents of a SYCL integration header file produced by a SYCL -// device compiler and used by SYCL host compiler (via forced inclusion into -// compiled SYCL source): -// - SYCL kernel names -// - SYCL kernel parameters and offsets of corresponding actual arguments -class SYCLIntegrationHeader { -public: - // Kind of kernel's parameters as captured by the compiler in the - // kernel lambda or function object - enum kernel_param_kind_t { - kind_first, - kind_accessor = kind_first, - kind_std_layout, - kind_sampler, - kind_pointer, - kind_specialization_constants_buffer, - kind_stream, - kind_last = kind_stream - }; - -public: - SYCLIntegrationHeader(Sema &S); - - /// Emits contents of the header into given stream. - void emit(raw_ostream &Out); - - /// Emits contents of the header into a file with given name. - /// Returns true/false on success/failure. - bool emit(StringRef MainSrc); - - /// Signals that subsequent parameter descriptor additions will go to - /// the kernel with given name. Starts new kernel invocation descriptor. - void startKernel(const FunctionDecl *SyclKernel, QualType KernelNameType, - SourceLocation Loc, bool IsESIMD, bool IsUnnamedKernel, - int64_t ObjSize); - - /// Adds a kernel parameter descriptor to current kernel invocation - /// descriptor. - void addParamDesc(kernel_param_kind_t Kind, int Info, unsigned Offset); - - /// Signals that addition of parameter descriptors to current kernel - /// invocation descriptor has finished. - void endKernel(); - - /// Registers a specialization constant to emit info for it into the header. - void addSpecConstant(StringRef IDName, QualType IDType); - - /// Update the names of a kernel description based on its SyclKernel. - void updateKernelNames(const FunctionDecl *SyclKernel, StringRef Name, - StringRef StableName) { - auto Itr = llvm::find_if(KernelDescs, [SyclKernel](const KernelDesc &KD) { - return KD.SyclKernel == SyclKernel; - }); - - assert(Itr != KernelDescs.end() && "Unknown kernel description"); - Itr->updateKernelNames(Name, StableName); - } - - /// Signals that emission of __sycl_device_global_registration type and - /// declaration of variable __sycl_device_global_registrar of this type in - /// integration header is required. - void addDeviceGlobalRegistration() { - NeedToEmitDeviceGlobalRegistration = true; - } - - /// Signals that emission of __sycl_host_pipe_registration type and - /// declaration of variable __sycl_host_pipe_registrar of this type in - /// integration header is required. - void addHostPipeRegistration() { - NeedToEmitHostPipeRegistration = true; - } - -private: - // Kernel actual parameter descriptor. - struct KernelParamDesc { - // Represents a parameter kind. - kernel_param_kind_t Kind = kind_last; - // If Kind is kind_scalar or kind_struct, then - // denotes parameter size in bytes (includes padding for structs) - // If Kind is kind_accessor - // denotes access target; possible access targets are defined in - // access/access.hpp - int Info = 0; - // Offset of the captured parameter value in the lambda or function object. - unsigned Offset = 0; - - KernelParamDesc() = default; - }; - - // Kernel invocation descriptor - struct KernelDesc { - /// sycl_kernel function associated with this kernel. - const FunctionDecl *SyclKernel; - - /// Kernel name. - std::string Name; - - /// Kernel name type. - QualType NameType; - - /// Kernel name with stable lambda name mangling - std::string StableName; - - SourceLocation KernelLocation; - - /// Whether this kernel is an ESIMD one. - bool IsESIMDKernel; - - /// Descriptor of kernel actual parameters. - SmallVector Params; - - // If we are in unnamed kernel/lambda mode AND this is one that the user - // hasn't provided an explicit name for. - bool IsUnnamedKernel; - - /// Size of the kernel object. - int64_t ObjSize = 0; - - KernelDesc(const FunctionDecl *SyclKernel, QualType NameType, - SourceLocation KernelLoc, bool IsESIMD, bool IsUnnamedKernel, - int64_t ObjSize) - : SyclKernel(SyclKernel), NameType(NameType), KernelLocation(KernelLoc), - IsESIMDKernel(IsESIMD), IsUnnamedKernel(IsUnnamedKernel), - ObjSize(ObjSize) {} - - void updateKernelNames(StringRef Name, StringRef StableName) { - this->Name = Name.str(); - this->StableName = StableName.str(); - } - }; - - /// Returns the latest invocation descriptor started by - /// SYCLIntegrationHeader::startKernel - KernelDesc *getCurKernelDesc() { - return KernelDescs.size() > 0 ? &KernelDescs[KernelDescs.size() - 1] - : nullptr; - } - -private: - /// Keeps invocation descriptors for each kernel invocation started by - /// SYCLIntegrationHeader::startKernel - SmallVector KernelDescs; - - using SpecConstID = std::pair; - - /// Keeps specialization constants met in the translation unit. Maps spec - /// constant's ID type to generated unique name. Duplicates are removed at - /// integration header emission time. - llvm::SmallVector SpecConsts; - - Sema &S; - - /// Keeps track of whether declaration of __sycl_device_global_registration - /// type and __sycl_device_global_registrar variable are required to emit. - bool NeedToEmitDeviceGlobalRegistration = false; - - /// Keeps track of whether declaration of __sycl_host_pipe_registration - /// type and __sycl_host_pipe_registrar variable are required to emit. - bool NeedToEmitHostPipeRegistration = false; -}; - -class SYCLIntegrationFooter { -public: - SYCLIntegrationFooter(Sema &S) : S(S) {} - bool emit(StringRef MainSrc); - void addVarDecl(const VarDecl *VD); - -private: - bool emit(raw_ostream &O); - Sema &S; - llvm::SmallVector GlobalVars; - void emitSpecIDName(raw_ostream &O, const VarDecl *VD); -}; - /// Tracks expected type during expression parsing, for use in code completion. /// The type is tied to a particular token, all functions that update or consume /// the type take a start location of the token they are looking at as a @@ -655,7 +469,6 @@ class Sema final : public SemaBase { // 37. Name Lookup for RISC-V Vector Intrinsic (SemaRISCVVectorLookup.cpp) // 38. CUDA (SemaCUDA.cpp) // 39. OpenMP Directives and Clauses (SemaOpenMP.cpp) - // 40. SYCL Constructs (SemaSYCL.cpp) /// \name Semantic Analysis /// Implementations are in Sema.cpp @@ -1162,6 +975,11 @@ class Sema final : public SemaBase { return *OpenACCPtr; } + SemaSYCL &SYCL() { + assert(SYCLPtr); + return *SYCLPtr; + } + protected: friend class Parser; friend class InitializationSequence; @@ -1194,6 +1012,7 @@ class Sema final : public SemaBase { std::unique_ptr HLSLPtr; std::unique_ptr OpenACCPtr; + std::unique_ptr SYCLPtr; ///@} @@ -5795,8 +5614,7 @@ class Sema final : public SemaBase { ExprResult BuildDeclarationNameExpr(const CXXScopeSpec &SS, LookupResult &R, bool NeedsADL, - bool AcceptInvalidDecl = false, - bool NeedUnresolved = false); + bool AcceptInvalidDecl = false); ExprResult BuildDeclarationNameExpr( const CXXScopeSpec &SS, const DeclarationNameInfo &NameInfo, NamedDecl *D, NamedDecl *FoundD = nullptr, @@ -5812,15 +5630,6 @@ class Sema final : public SemaBase { ExprResult ActOnPredefinedExpr(SourceLocation Loc, tok::TokenKind Kind); ExprResult ActOnIntegerConstant(SourceLocation Loc, uint64_t Val); - ExprResult BuildSYCLUniqueStableNameExpr(SourceLocation OpLoc, - SourceLocation LParen, - SourceLocation RParen, - TypeSourceInfo *TSI); - ExprResult ActOnSYCLUniqueStableNameExpr(SourceLocation OpLoc, - SourceLocation LParen, - SourceLocation RParen, - ParsedType ParsedTy); - bool CheckLoopHintExpr(Expr *E, SourceLocation Loc); ExprResult ActOnNumericConstant(const Token &Tok, Scope *UDLScope = nullptr); @@ -6956,10 +6765,7 @@ class Sema final : public SemaBase { SourceLocation RParenLoc); //// ActOnCXXThis - Parse 'this' pointer. - ExprResult ActOnCXXThis(SourceLocation Loc); - - /// Check whether the type of 'this' is valid in the current context. - bool CheckCXXThisType(SourceLocation Loc, QualType Type); + ExprResult ActOnCXXThis(SourceLocation loc); /// Build a CXXThisExpr and mark it referenced in the current context. Expr *BuildCXXThisExpr(SourceLocation Loc, QualType Type, bool IsImplicit); @@ -14907,177 +14713,6 @@ class Sema final : public SemaBase { const DeclarationNameInfo &DirName, OpenMPDirectiveKind CancelRegion); - ///@} - // - // - // ------------------------------------------------------------------------- - // - // - - /// \name SYCL Constructs - /// Implementations are in SemaSYCL.cpp - ///@{ - -private: - - void CheckSYCLKernelCall(FunctionDecl *CallerFunc, - ArrayRef Args); - - // We store SYCL Kernels here and handle separately -- which is a hack. - // FIXME: It would be best to refactor this. - llvm::SetVector SyclDeviceDecls; - // SYCL integration header instance for current compilation unit this Sema - // is associated with. - std::unique_ptr SyclIntHeader; - std::unique_ptr SyclIntFooter; - - // We need to store the list of the sycl_kernel functions and their associated - // generated OpenCL Kernels so we can go back and re-name these after the - // fact. - llvm::SmallVector> - SyclKernelsToOpenCLKernels; - - // Used to suppress diagnostics during kernel construction, since these were - // already emitted earlier. Diagnosing during Kernel emissions also skips the - // useful notes that shows where the kernel was called. - bool DiagnosingSYCLKernel = false; - -public: - void addSyclOpenCLKernel(const FunctionDecl *SyclKernel, - FunctionDecl *OpenCLKernel) { - SyclKernelsToOpenCLKernels.emplace_back(SyclKernel, OpenCLKernel); - } - - void addSyclDeviceDecl(Decl *d) { SyclDeviceDecls.insert(d); } - llvm::SetVector &syclDeviceDecls() { return SyclDeviceDecls; } - - /// Lazily creates and returns SYCL integration header instance. - SYCLIntegrationHeader &getSyclIntegrationHeader() { - if (SyclIntHeader == nullptr) - SyclIntHeader = std::make_unique(*this); - return *SyclIntHeader.get(); - } - - SYCLIntegrationFooter &getSyclIntegrationFooter() { - if (SyclIntFooter == nullptr) - SyclIntFooter = std::make_unique(*this); - return *SyclIntFooter.get(); - } - - void addSyclVarDecl(VarDecl *VD) { - if (LangOpts.SYCLIsDevice && !LangOpts.SYCLIntFooter.empty()) - getSyclIntegrationFooter().addVarDecl(VD); - } - - enum SYCLRestrictKind { - KernelGlobalVariable, - KernelRTTI, - KernelNonConstStaticDataVariable, - KernelCallVirtualFunction, - KernelUseExceptions, - KernelCallRecursiveFunction, - KernelCallFunctionPointer, - KernelAllocateStorage, - KernelUseAssembly, - KernelCallDllimportFunction, - KernelCallVariadicFunction, - KernelCallUndefinedFunction, - KernelConstStaticVariable - }; - - bool isDeclAllowedInSYCLDeviceCode(const Decl *D); - void checkSYCLDeviceVarDecl(VarDecl *Var); - void copySYCLKernelAttrs(CXXMethodDecl *CallOperator); - void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); - void SetSYCLKernelNames(); - void MarkDevices(); - - /// Get the number of fields or captures within the parsed type. - ExprResult ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT); - ExprResult BuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, - QualType SourceTy); - - /// Get a value based on the type of the given field number so that callers - /// can wrap it in a decltype() to get the actual type of the field. - ExprResult ActOnSYCLBuiltinFieldTypeExpr(ParsedType PT, Expr *Idx); - ExprResult BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, - QualType SourceTy, Expr *Idx); - - /// Get the number of base classes within the parsed type. - ExprResult ActOnSYCLBuiltinNumBasesExpr(ParsedType PT); - ExprResult BuildSYCLBuiltinNumBasesExpr(SourceLocation Loc, - QualType SourceTy); - - /// Get a value based on the type of the given base number so that callers - /// can wrap it in a decltype() to get the actual type of the base class. - ExprResult ActOnSYCLBuiltinBaseTypeExpr(ParsedType PT, Expr *Idx); - ExprResult BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, QualType SourceTy, - Expr *Idx); - - bool checkAllowedSYCLInitializer(VarDecl *VD); - - /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current - /// context is "used as device code". - /// - /// - If CurLexicalContext is a kernel function or it is known that the - /// function will be emitted for the device, emits the diagnostics - /// immediately. - /// - If CurLexicalContext is a function and we are compiling - /// for the device, but we don't know that this function will be codegen'ed - /// for device yet, creates a diagnostic which is emitted if and when we - /// realize that the function will be codegen'ed. - /// - /// Example usage: - /// - /// Diagnose __float128 type usage only from SYCL device code if the current - /// target doesn't support it - /// if (!S.Context.getTargetInfo().hasFloat128Type() && - /// S.getLangOpts().SYCLIsDevice) - /// SYCLDiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128"; - SemaDiagnosticBuilder SYCLDiagIfDeviceCode( - SourceLocation Loc, unsigned DiagID, - DeviceDiagnosticReason Reason = DeviceDiagnosticReason::Sycl | - DeviceDiagnosticReason::Esimd); - - void deepTypeCheckForSYCLDevice(SourceLocation UsedAt, - llvm::DenseSet Visited, - ValueDecl *DeclToCheck); - - /// Finishes analysis of the deferred functions calls that may be not - /// properly declared for device compilation. - void finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, - const FunctionDecl *Callee, - SourceLocation Loc, - DeviceDiagnosticReason Reason); - - /// Tells whether given variable is a SYCL explicit SIMD extension's "private - /// global" variable - global variable in the private address space. - bool isSYCLEsimdPrivateGlobal(VarDecl *VDecl) { - return getLangOpts().SYCLIsDevice && VDecl->hasAttr() && - VDecl->hasGlobalStorage() && - (VDecl->getType().getAddressSpace() == LangAS::sycl_private); - } - - template - static bool isTypeDecoratedWithDeclAttribute(QualType Ty) { - const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); - if (!RecTy) - return false; - - if (RecTy->hasAttr()) - return true; - - if (auto *CTSD = dyn_cast(RecTy)) { - ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); - if (CXXRecordDecl *RD = Template->getTemplatedDecl()) - return RD->hasAttr(); - } - return false; - } - - /// Check whether \p Ty corresponds to a SYCL type of name \p TypeName. - static bool isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName); - ///@} }; diff --git a/clang/include/clang/Sema/SemaSYCL.h b/clang/include/clang/Sema/SemaSYCL.h new file mode 100644 index 0000000000000..9b6d883996bcb --- /dev/null +++ b/clang/include/clang/Sema/SemaSYCL.h @@ -0,0 +1,402 @@ +//===----- SemaSYCL.h ------- Semantic Analysis for SYCL constructs -------===// +// +// 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 +// +//===----------------------------------------------------------------------===// +/// \file +/// This file declares semantic analysis for SYCL constructs. +/// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_SEMA_SEMASYCL_H +#define LLVM_CLANG_SEMA_SEMASYCL_H + +#include "clang/AST/Attr.h" +#include "clang/AST/Decl.h" +#include "clang/AST/Type.h" +#include "clang/Basic/SourceLocation.h" +#include "clang/Sema/Ownership.h" +#include "clang/Sema/SemaBase.h" +#include "llvm/ADT/DenseSet.h" +#include "llvm/ADT/SetVector.h" + +namespace clang { + +class CXXMethodDecl; +class MangleContext; +class SemaSYCL; + +// TODO SYCL Integration header approach relies on an assumption that kernel +// lambda objects created by the host compiler and any of the device compilers +// will be identical wrt to field types, order and offsets. Some verification +// mechanism should be developed to enforce that. + +// TODO FIXME SYCL Support for SYCL in FE should be refactored: +// - kernel identification and generation should be made a separate pass over +// AST. RecursiveASTVisitor + VisitFunctionTemplateDecl + +// FunctionTemplateDecl::getSpecializations() mechanism could be used for that. +// - All SYCL stuff on Sema level should be encapsulated into a single Sema +// field +// - Move SYCL stuff into a separate header + +// Represents contents of a SYCL integration header file produced by a SYCL +// device compiler and used by SYCL host compiler (via forced inclusion into +// compiled SYCL source): +// - SYCL kernel names +// - SYCL kernel parameters and offsets of corresponding actual arguments +class SYCLIntegrationHeader { +public: + // Kind of kernel's parameters as captured by the compiler in the + // kernel lambda or function object + enum kernel_param_kind_t { + kind_first, + kind_accessor = kind_first, + kind_std_layout, + kind_sampler, + kind_pointer, + kind_specialization_constants_buffer, + kind_stream, + kind_last = kind_stream + }; + +public: + SYCLIntegrationHeader(SemaSYCL &S); + + /// Emits contents of the header into given stream. + void emit(raw_ostream &Out); + + /// Emits contents of the header into a file with given name. + /// Returns true/false on success/failure. + bool emit(StringRef MainSrc); + + /// Signals that subsequent parameter descriptor additions will go to + /// the kernel with given name. Starts new kernel invocation descriptor. + void startKernel(const FunctionDecl *SyclKernel, QualType KernelNameType, + SourceLocation Loc, bool IsESIMD, bool IsUnnamedKernel, + int64_t ObjSize); + + /// Adds a kernel parameter descriptor to current kernel invocation + /// descriptor. + void addParamDesc(kernel_param_kind_t Kind, int Info, unsigned Offset); + + /// Signals that addition of parameter descriptors to current kernel + /// invocation descriptor has finished. + void endKernel(); + + /// Registers a specialization constant to emit info for it into the header. + void addSpecConstant(StringRef IDName, QualType IDType); + + /// Update the names of a kernel description based on its SyclKernel. + void updateKernelNames(const FunctionDecl *SyclKernel, StringRef Name, + StringRef StableName) { + auto Itr = llvm::find_if(KernelDescs, [SyclKernel](const KernelDesc &KD) { + return KD.SyclKernel == SyclKernel; + }); + + assert(Itr != KernelDescs.end() && "Unknown kernel description"); + Itr->updateKernelNames(Name, StableName); + } + + /// Signals that emission of __sycl_device_global_registration type and + /// declaration of variable __sycl_device_global_registrar of this type in + /// integration header is required. + void addDeviceGlobalRegistration() { + NeedToEmitDeviceGlobalRegistration = true; + } + + /// Signals that emission of __sycl_host_pipe_registration type and + /// declaration of variable __sycl_host_pipe_registrar of this type in + /// integration header is required. + void addHostPipeRegistration() { + NeedToEmitHostPipeRegistration = true; + } + +private: + // Kernel actual parameter descriptor. + struct KernelParamDesc { + // Represents a parameter kind. + kernel_param_kind_t Kind = kind_last; + // If Kind is kind_scalar or kind_struct, then + // denotes parameter size in bytes (includes padding for structs) + // If Kind is kind_accessor + // denotes access target; possible access targets are defined in + // access/access.hpp + int Info = 0; + // Offset of the captured parameter value in the lambda or function object. + unsigned Offset = 0; + + KernelParamDesc() = default; + }; + + // Kernel invocation descriptor + struct KernelDesc { + /// sycl_kernel function associated with this kernel. + const FunctionDecl *SyclKernel; + + /// Kernel name. + std::string Name; + + /// Kernel name type. + QualType NameType; + + /// Kernel name with stable lambda name mangling + std::string StableName; + + SourceLocation KernelLocation; + + /// Whether this kernel is an ESIMD one. + bool IsESIMDKernel; + + /// Descriptor of kernel actual parameters. + SmallVector Params; + + // If we are in unnamed kernel/lambda mode AND this is one that the user + // hasn't provided an explicit name for. + bool IsUnnamedKernel; + + /// Size of the kernel object. + int64_t ObjSize = 0; + + KernelDesc(const FunctionDecl *SyclKernel, QualType NameType, + SourceLocation KernelLoc, bool IsESIMD, bool IsUnnamedKernel, + int64_t ObjSize) + : SyclKernel(SyclKernel), NameType(NameType), KernelLocation(KernelLoc), + IsESIMDKernel(IsESIMD), IsUnnamedKernel(IsUnnamedKernel), + ObjSize(ObjSize) {} + + void updateKernelNames(StringRef Name, StringRef StableName) { + this->Name = Name.str(); + this->StableName = StableName.str(); + } + }; + + /// Returns the latest invocation descriptor started by + /// SYCLIntegrationHeader::startKernel + KernelDesc *getCurKernelDesc() { + return KernelDescs.size() > 0 ? &KernelDescs[KernelDescs.size() - 1] + : nullptr; + } + +private: + /// Keeps invocation descriptors for each kernel invocation started by + /// SYCLIntegrationHeader::startKernel + SmallVector KernelDescs; + + using SpecConstID = std::pair; + + /// Keeps specialization constants met in the translation unit. Maps spec + /// constant's ID type to generated unique name. Duplicates are removed at + /// integration header emission time. + llvm::SmallVector SpecConsts; + + SemaSYCL &S; + + /// Keeps track of whether declaration of __sycl_device_global_registration + /// type and __sycl_device_global_registrar variable are required to emit. + bool NeedToEmitDeviceGlobalRegistration = false; + + /// Keeps track of whether declaration of __sycl_host_pipe_registration + /// type and __sycl_host_pipe_registrar variable are required to emit. + bool NeedToEmitHostPipeRegistration = false; +}; + +class SYCLIntegrationFooter { +public: + SYCLIntegrationFooter(SemaSYCL &S) : S(S) {} + bool emit(StringRef MainSrc); + void addVarDecl(const VarDecl *VD); + +private: + bool emit(raw_ostream &O); + SemaSYCL &S; + llvm::SmallVector GlobalVars; + void emitSpecIDName(raw_ostream &O, const VarDecl *VD); +}; + +class SemaSYCL : public SemaBase { +private: + + // We store SYCL Kernels here and handle separately -- which is a hack. + // FIXME: It would be best to refactor this. + llvm::SetVector SyclDeviceDecls; + // SYCL integration header instance for current compilation unit this Sema + // is associated with. + std::unique_ptr SyclIntHeader; + std::unique_ptr SyclIntFooter; + + // We need to store the list of the sycl_kernel functions and their associated + // generated OpenCL Kernels so we can go back and re-name these after the + // fact. + llvm::SmallVector> + SyclKernelsToOpenCLKernels; + + // Used to suppress diagnostics during kernel construction, since these were + // already emitted earlier. Diagnosing during Kernel emissions also skips the + // useful notes that shows where the kernel was called. + bool DiagnosingSYCLKernel = false; + +public: + SemaSYCL(Sema &S); + + void CheckSYCLKernelCall(FunctionDecl *CallerFunc, + ArrayRef Args); + + /// Creates a SemaDiagnosticBuilder that emits the diagnostic if the current + /// context is "used as device code". + /// + /// - If CurLexicalContext is a kernel function or it is known that the + /// function will be emitted for the device, emits the diagnostics + /// immediately. + /// - If CurLexicalContext is a function and we are compiling + /// for the device, but we don't know that this function will be codegen'ed + /// for device yet, creates a diagnostic which is emitted if and when we + /// realize that the function will be codegen'ed. + /// + /// Example usage: + /// + /// Diagnose __float128 type usage only from SYCL device code if the current + /// target doesn't support it + /// if (!S.Context.getTargetInfo().hasFloat128Type() && + /// S.getLangOpts().SYCLIsDevice) + /// DiagIfDeviceCode(Loc, diag::err_type_unsupported) << "__float128"; + SemaDiagnosticBuilder DiagIfDeviceCode( + SourceLocation Loc, unsigned DiagID, + DeviceDiagnosticReason Reason = DeviceDiagnosticReason::Sycl | + DeviceDiagnosticReason::Esimd); + + void deepTypeCheckForDevice(SourceLocation UsedAt, + llvm::DenseSet Visited, + ValueDecl *DeclToCheck); + + void addSyclOpenCLKernel(const FunctionDecl *SyclKernel, + FunctionDecl *OpenCLKernel) { + SyclKernelsToOpenCLKernels.emplace_back(SyclKernel, OpenCLKernel); + } + + void addSyclDeviceDecl(Decl *d) { SyclDeviceDecls.insert(d); } + llvm::SetVector &syclDeviceDecls() { return SyclDeviceDecls; } + + /// Lazily creates and returns SYCL integration header instance. + SYCLIntegrationHeader &getSyclIntegrationHeader() { + if (SyclIntHeader == nullptr) + SyclIntHeader = std::make_unique(*this); + return *SyclIntHeader.get(); + } + + SYCLIntegrationFooter &getSyclIntegrationFooter() { + if (SyclIntFooter == nullptr) + SyclIntFooter = std::make_unique(*this); + return *SyclIntFooter.get(); + } + + void addSyclVarDecl(VarDecl *VD) { + if (getLangOpts().SYCLIsDevice && !getLangOpts().SYCLIntFooter.empty()) + getSyclIntegrationFooter().addVarDecl(VD); + } + + bool hasSyclIntegrationHeader() { return SyclIntHeader != nullptr; } + bool hasSyclIntegrationFooter() { return SyclIntFooter != nullptr; } + + enum SYCLRestrictKind { + KernelGlobalVariable, + KernelRTTI, + KernelNonConstStaticDataVariable, + KernelCallVirtualFunction, + KernelUseExceptions, + KernelCallRecursiveFunction, + KernelCallFunctionPointer, + KernelAllocateStorage, + KernelUseAssembly, + KernelCallDllimportFunction, + KernelCallVariadicFunction, + KernelCallUndefinedFunction, + KernelConstStaticVariable + }; + + bool isDeclAllowedInSYCLDeviceCode(const Decl *D); + void checkSYCLDeviceVarDecl(VarDecl *Var); + void copySYCLKernelAttrs(CXXMethodDecl *CallOperator); + void ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC); + void SetSYCLKernelNames(); + void MarkDevices(); + + /// Get the number of fields or captures within the parsed type. + ExprResult ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT); + ExprResult BuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, + QualType SourceTy); + + /// Get a value based on the type of the given field number so that callers + /// can wrap it in a decltype() to get the actual type of the field. + ExprResult ActOnSYCLBuiltinFieldTypeExpr(ParsedType PT, Expr *Idx); + ExprResult BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, + QualType SourceTy, Expr *Idx); + + /// Get the number of base classes within the parsed type. + ExprResult ActOnSYCLBuiltinNumBasesExpr(ParsedType PT); + ExprResult BuildSYCLBuiltinNumBasesExpr(SourceLocation Loc, + QualType SourceTy); + + /// Get a value based on the type of the given base number so that callers + /// can wrap it in a decltype() to get the actual type of the base class. + ExprResult ActOnSYCLBuiltinBaseTypeExpr(ParsedType PT, Expr *Idx); + ExprResult BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, QualType SourceTy, + Expr *Idx); + + bool checkAllowedSYCLInitializer(VarDecl *VD); + + /// Finishes analysis of the deferred functions calls that may be not + /// properly declared for device compilation. + void finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, + const FunctionDecl *Callee, + SourceLocation Loc, + DeviceDiagnosticReason Reason); + + /// Tells whether given variable is a SYCL explicit SIMD extension's "private + /// global" variable - global variable in the private address space. + bool isSYCLEsimdPrivateGlobal(VarDecl *VDecl) { + return getLangOpts().SYCLIsDevice && VDecl->hasAttr() && + VDecl->hasGlobalStorage() && + (VDecl->getType().getAddressSpace() == LangAS::sycl_private); + } + + template + static bool isTypeDecoratedWithDeclAttribute(QualType Ty) { + const CXXRecordDecl *RecTy = Ty->getAsCXXRecordDecl(); + if (!RecTy) + return false; + + if (RecTy->hasAttr()) + return true; + + if (auto *CTSD = dyn_cast(RecTy)) { + ClassTemplateDecl *Template = CTSD->getSpecializedTemplate(); + if (CXXRecordDecl *RD = Template->getTemplatedDecl()) + return RD->hasAttr(); + } + return false; + } + + /// Check whether \p Ty corresponds to a SYCL type of name \p TypeName. + static bool isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName); + + ExprResult BuildUniqueStableIdExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, Expr *E); + ExprResult ActOnUniqueStableIdExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, Expr *E); + ExprResult BuildUniqueStableNameExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, + TypeSourceInfo *TSI); + ExprResult ActOnUniqueStableNameExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, + ParsedType ParsedTy); +}; + +} // namespace clang + +#endif // LLVM_CLANG_SEMA_SEMASYCL_H diff --git a/clang/lib/AST/TypePrinter.cpp b/clang/lib/AST/TypePrinter.cpp index a66aef6323a37..55308d153724a 100644 --- a/clang/lib/AST/TypePrinter.cpp +++ b/clang/lib/AST/TypePrinter.cpp @@ -1222,10 +1222,13 @@ void TypePrinter::printDecltypeBefore(const DecltypeType *T, raw_ostream &OS) { void TypePrinter::printPackIndexingBefore(const PackIndexingType *T, raw_ostream &OS) { - if (T->hasSelectedType()) + if (T->hasSelectedType()) { OS << T->getSelectedType(); - else - OS << T->getPattern() << "...[" << T->getIndexExpr() << "]"; + } else { + OS << T->getPattern() << "...["; + T->getIndexExpr()->printPretty(OS, nullptr, Policy); + OS << "]"; + } spaceBeforePlaceHolder(OS); } diff --git a/clang/lib/CodeGen/CodeGenModule.cpp b/clang/lib/CodeGen/CodeGenModule.cpp index 5c4feb4994431..145b76507b162 100644 --- a/clang/lib/CodeGen/CodeGenModule.cpp +++ b/clang/lib/CodeGen/CodeGenModule.cpp @@ -51,6 +51,7 @@ #include "clang/CodeGen/ConstantInitBuilder.h" #include "clang/Frontend/FrontendDiagnostic.h" #include "clang/Sema/Sema.h" +#include "clang/Sema/SemaSYCL.h" #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/StringExtras.h" #include "llvm/ADT/StringSwitch.h" @@ -6254,7 +6255,7 @@ CodeGenModule::getLLVMLinkageForDeclarator(const DeclaratorDecl *D, // is only one translation unit and can so mark them internal. if (getLangOpts().SYCLIsDevice && !D->hasAttr() && !D->hasAttr() && - !Sema::isTypeDecoratedWithDeclAttribute( + !SemaSYCL::isTypeDecoratedWithDeclAttribute( D->getType())) return getLangOpts().GPURelocatableDeviceCode ? llvm::Function::LinkOnceODRLinkage diff --git a/clang/lib/Parse/ParseAST.cpp b/clang/lib/Parse/ParseAST.cpp index 4e47ea5ea364f..dc8df8e8f3f03 100644 --- a/clang/lib/Parse/ParseAST.cpp +++ b/clang/lib/Parse/ParseAST.cpp @@ -21,6 +21,7 @@ #include "clang/Sema/EnterExpressionEvaluationContext.h" #include "clang/Sema/Sema.h" #include "clang/Sema/SemaConsumer.h" +#include "clang/Sema/SemaSYCL.h" #include "clang/Sema/TemplateInstCallback.h" #include "llvm/Support/CrashRecoveryContext.h" #include "llvm/Support/TimeProfiler.h" @@ -174,7 +175,7 @@ void clang::ParseAST(Sema &S, bool PrintStats, bool SkipFunctionBodies) { Consumer->HandleTopLevelDecl(DeclGroupRef(D)); if (S.getLangOpts().SYCLIsDevice) { - for (Decl *D : S.syclDeviceDecls()) { + for (Decl *D : S.SYCL().syclDeviceDecls()) { Consumer->HandleTopLevelDecl(DeclGroupRef(D)); } } diff --git a/clang/lib/Parse/ParseExpr.cpp b/clang/lib/Parse/ParseExpr.cpp index 7164b3e34b028..7b15590170fd4 100644 --- a/clang/lib/Parse/ParseExpr.cpp +++ b/clang/lib/Parse/ParseExpr.cpp @@ -30,6 +30,7 @@ #include "clang/Sema/EnterExpressionEvaluationContext.h" #include "clang/Sema/ParsedTemplate.h" #include "clang/Sema/Scope.h" +#include "clang/Sema/SemaSYCL.h" #include "clang/Sema/TypoCorrection.h" #include "llvm/ADT/SmallVector.h" #include @@ -1949,8 +1950,8 @@ ExprResult Parser::ParseSYCLBuiltinNum() { T.consumeClose(); if (IsNumFields) - return Actions.ActOnSYCLBuiltinNumFieldsExpr(TR.get()); - return Actions.ActOnSYCLBuiltinNumBasesExpr(TR.get()); + return Actions.SYCL().ActOnSYCLBuiltinNumFieldsExpr(TR.get()); + return Actions.SYCL().ActOnSYCLBuiltinNumBasesExpr(TR.get()); } /// __builtin_field_type '(' type-id ',' integer-constant ')' or @@ -1981,8 +1982,8 @@ ExprResult Parser::ParseSYCLBuiltinType() { T.consumeClose(); if (IsFieldType) - return Actions.ActOnSYCLBuiltinFieldTypeExpr(TR.get(), IdxRes.get()); - return Actions.ActOnSYCLBuiltinBaseTypeExpr(TR.get(), IdxRes.get()); + return Actions.SYCL().ActOnSYCLBuiltinFieldTypeExpr(TR.get(), IdxRes.get()); + return Actions.SYCL().ActOnSYCLBuiltinBaseTypeExpr(TR.get(), IdxRes.get()); } /// Once the leading part of a postfix-expression is parsed, this @@ -2559,8 +2560,8 @@ ExprResult Parser::ParseSYCLUniqueStableNameExpression() { if (T.consumeClose()) return ExprError(); - return Actions.ActOnSYCLUniqueStableNameExpr(OpLoc, T.getOpenLocation(), - T.getCloseLocation(), Ty.get()); + return Actions.SYCL().ActOnUniqueStableNameExpr( + OpLoc, T.getOpenLocation(), T.getCloseLocation(), Ty.get()); } // Parse a __builtin_sycl_unique_stable_id expression. Accepts an expression, @@ -2590,7 +2591,7 @@ ExprResult Parser::ParseSYCLUniqueStableIdExpression() { if (T.consumeClose()) return ExprError(); - return Actions.ActOnSYCLUniqueStableIdExpr( + return Actions.SYCL().ActOnUniqueStableIdExpr( OpLoc, T.getOpenLocation(), T.getCloseLocation(), VarExpr.get()); } diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index 09b56fddea882..901ca0aea01d8 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -45,6 +45,7 @@ #include "clang/Sema/SemaHLSL.h" #include "clang/Sema/SemaInternal.h" #include "clang/Sema/SemaOpenACC.h" +#include "clang/Sema/SemaSYCL.h" #include "clang/Sema/TemplateDeduction.h" #include "clang/Sema/TemplateInstCallback.h" #include "clang/Sema/TypoCorrection.h" @@ -201,6 +202,7 @@ Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer, CurScope(nullptr), Ident_super(nullptr), HLSLPtr(std::make_unique(*this)), OpenACCPtr(std::make_unique(*this)), + SYCLPtr(std::make_unique(*this)), MSPointerToMemberRepresentationMethod( LangOpts.getMSPointerToMemberRepresentationMethod()), MSStructPragmaOn(false), VtorDispStack(LangOpts.getVtorDispMode()), @@ -224,8 +226,7 @@ Sema::Sema(Preprocessor &pp, ASTContext &ctxt, ASTConsumer &consumer, ValueWithBytesObjCTypeMethod(nullptr), NSArrayDecl(nullptr), ArrayWithObjectsMethod(nullptr), NSDictionaryDecl(nullptr), DictionaryWithObjectsMethod(nullptr), CodeCompleter(CodeCompleter), - VarDataSharingAttributesStack(nullptr), - SyclIntHeader(nullptr), SyclIntFooter(nullptr) { + VarDataSharingAttributesStack(nullptr) { assert(pp.TUKind == TUKind); TUScope = nullptr; @@ -1125,16 +1126,16 @@ void Sema::ActOnEndOfTranslationUnitFragment(TUFragmentKind Kind) { if (getLangOpts().SYCLIsDevice) { // Set the names of the kernels, now that the names have settled down. This // needs to happen before we generate the integration headers. - SetSYCLKernelNames(); + SYCL().SetSYCLKernelNames(); // Make sure that the footer is emitted before header, since only after the // footer is emitted is it known that translation unit contains device // global variables. - if (SyclIntFooter != nullptr) - SyclIntFooter->emit(getLangOpts().SYCLIntFooter); + if (SYCL().hasSyclIntegrationFooter()) + SYCL().getSyclIntegrationFooter().emit(getLangOpts().SYCLIntFooter); // Emit SYCL integration header for current translation unit if needed - if (SyclIntHeader != nullptr) - SyclIntHeader->emit(getLangOpts().SYCLIntHeader); - MarkDevices(); + if (SYCL().hasSyclIntegrationHeader()) + SYCL().getSyclIntegrationHeader().emit(getLangOpts().SYCLIntHeader); + SYCL().MarkDevices(); } emitDeferredDiags(); @@ -1760,15 +1761,16 @@ class DeferredDiagnosticsEmitter void visitUsedDecl(SourceLocation Loc, Decl *D) { if (S.LangOpts.SYCLIsDevice && ShouldEmitRootNode) { if (auto *VD = dyn_cast(D)) { - if (!S.checkAllowedSYCLInitializer(VD) && - !S.isTypeDecoratedWithDeclAttribute( - VD->getType())) { + if (!S.SYCL().checkAllowedSYCLInitializer(VD) && + !S.SYCL() + .isTypeDecoratedWithDeclAttribute< + SYCLGlobalVariableAllowedAttr>(VD->getType())) { S.Diag(Loc, diag::err_sycl_restrict) - << Sema::KernelConstStaticVariable; + << SemaSYCL::KernelConstStaticVariable; return; } if (!VD->hasInit() && - S.isTypeDecoratedWithDeclAttribute( + S.SYCL().isTypeDecoratedWithDeclAttribute( VD->getType()) && !VD->hasAttr()) S.Diag(Loc, diag::err_sycl_external_global); @@ -1819,7 +1821,7 @@ class DeferredDiagnosticsEmitter S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc); // Finalize analysis of SYCL-specific constructs. if (Caller && S.LangOpts.SYCLIsDevice) - S.finalizeSYCLDelayedAnalysis(Caller, FD, Loc, RootReason); + S.SYCL().finalizeSYCLDelayedAnalysis(Caller, FD, Loc, RootReason); if (Caller) S.DeviceKnownEmittedFns[FD] = {Caller, Loc}; // Always emit deferred diagnostics for the direct users. This does not @@ -1994,7 +1996,7 @@ Sema::targetDiag(SourceLocation Loc, unsigned DiagID, const FunctionDecl *FD) { : diagIfOpenMPHostCode(Loc, DiagID, FD); if (getLangOpts().SYCLIsDevice) - return SYCLDiagIfDeviceCode(Loc, DiagID); + return SYCL().DiagIfDeviceCode(Loc, DiagID); if (getLangOpts().CUDA) return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID) @@ -2014,7 +2016,7 @@ void Sema::checkTypeSupport(QualType Ty, SourceLocation Loc, ValueDecl *D) { // constant byte size like zero length arrays. So, do a deep check for SYCL. if (D && LangOpts.SYCLIsDevice) { llvm::DenseSet Visited; - deepTypeCheckForSYCLDevice(Loc, Visited, D); + SYCL().deepTypeCheckForDevice(Loc, Visited, D); } Decl *C = cast(getCurLexicalContext()); diff --git a/clang/lib/Sema/SemaCast.cpp b/clang/lib/Sema/SemaCast.cpp index bc4441f09028f..1c479b85be2a3 100644 --- a/clang/lib/Sema/SemaCast.cpp +++ b/clang/lib/Sema/SemaCast.cpp @@ -24,6 +24,7 @@ #include "clang/Lex/Preprocessor.h" #include "clang/Sema/Initialization.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "llvm/ADT/SmallVector.h" #include "llvm/ADT/StringExtras.h" #include @@ -2634,8 +2635,8 @@ static TryCastResult TryAddressSpaceCast(Sema &Self, ExprResult &SrcExpr, Qualifiers SrcQ = SrcPointeeType.getQualifiers(); Qualifiers DestQ = DestPointeeType.getQualifiers(); if (!DestQ.isAddressSpaceSupersetOf(SrcQ) && OpRange.isValid()) { - Self.SYCLDiagIfDeviceCode(OpRange.getBegin(), - diag::warn_sycl_potentially_invalid_as_cast) + Self.SYCL().DiagIfDeviceCode(OpRange.getBegin(), + diag::warn_sycl_potentially_invalid_as_cast) << SrcType << DestType << OpRange; } } diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp index 8268d7b8d6c80..f34e8a3ed3a3d 100644 --- a/clang/lib/Sema/SemaChecking.cpp +++ b/clang/lib/Sema/SemaChecking.cpp @@ -62,6 +62,7 @@ #include "clang/Sema/ScopeInfo.h" #include "clang/Sema/Sema.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "llvm/ADT/APFloat.h" #include "llvm/ADT/APInt.h" #include "llvm/ADT/APSInt.h" @@ -3295,8 +3296,8 @@ Sema::CheckBuiltinFunctionCall(FunctionDecl *FDecl, unsigned BuiltinID, // Detect when host builtins are used in device code only if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(TheCall->getBeginLoc(), - diag::err_builtin_target_unsupported); + SYCL().DiagIfDeviceCode(TheCall->getBeginLoc(), + diag::err_builtin_target_unsupported); } else { if (CheckTSBuiltinFunctionCall(Context.getTargetInfo(), BuiltinID, TheCall)) @@ -7823,7 +7824,7 @@ bool Sema::CheckIntelSYCLAllocaBuiltinFunctionCall(unsigned, CallExpr *Call) { return true; Ty = Ty->getPointeeType(); return !(Ty.getQualifiers().empty() && - isSyclType(Ty, SYCLTypeAttr::kernel_handler)); + SemaSYCL::isSyclType(Ty, SYCLTypeAttr::kernel_handler)); }; if (CheckArg(FD->getParamDecl(0)->getType())) { Diag(Loc, diag::err_intel_sycl_alloca_wrong_arg) @@ -7835,7 +7836,7 @@ bool Sema::CheckIntelSYCLAllocaBuiltinFunctionCall(unsigned, CallExpr *Call) { // sycl::access::address_space::private_space, DecoratedAddress>`: // - `ET`: cv-unqualified trivial type constexpr auto CheckType = [](QualType RT, const ASTContext &Ctx) { - if (!isSyclType(RT, SYCLTypeAttr::multi_ptr)) + if (!SemaSYCL::isSyclType(RT, SYCLTypeAttr::multi_ptr)) return true; // Check element type const TemplateArgumentList &TAL = @@ -7860,7 +7861,7 @@ bool Sema::CheckIntelSYCLAllocaBuiltinFunctionCall(unsigned, CallExpr *Call) { if (Ty.isNull() || !Ty->isReferenceType()) return true; Ty = Ty->getPointeeType(); - if (!isSyclType(Ty, SYCLTypeAttr::specialization_id)) + if (!SemaSYCL::isSyclType(Ty, SYCLTypeAttr::specialization_id)) return true; const TemplateArgumentList &TAL = cast(Ty->getAsCXXRecordDecl()) @@ -7880,7 +7881,6 @@ bool Sema::CheckIntelSYCLAllocaBuiltinFunctionCall(unsigned, CallExpr *Call) { return false; } - /// Given a FunctionDecl's FormatAttr, attempts to populate the FomatStringInfo /// parameter with the FormatAttr's correct format_idx and firstDataArg. /// Returns true when the format fits the function and the FormatStringInfo has @@ -8361,13 +8361,13 @@ void Sema::checkCall(NamedDecl *FDecl, const FunctionProtoType *Proto, diagnoseArgDependentDiagnoseIfAttrs(FD, ThisArg, Args, Loc); if (FD && FD->hasAttr()) - CheckSYCLKernelCall(FD, Args); + SYCL().CheckSYCLKernelCall(FD, Args); // Diagnose variadic calls in SYCL. if (FD && FD->isVariadic() && getLangOpts().SYCLIsDevice && - !isUnevaluatedContext() && !isDeclAllowedInSYCLDeviceCode(FD)) - SYCLDiagIfDeviceCode(Loc, diag::err_sycl_restrict) - << Sema::KernelCallVariadicFunction; + !isUnevaluatedContext() && !SYCL().isDeclAllowedInSYCLDeviceCode(FD)) + SYCL().DiagIfDeviceCode(Loc, diag::err_sycl_restrict) + << SemaSYCL::KernelCallVariadicFunction; } /// CheckConstructorCall - Check a constructor call for correctness and safety @@ -16572,7 +16572,8 @@ static void CheckImplicitConversion(Sema &S, Expr *E, QualType T, S.Context.getFloatTypeSemantics(QualType(TargetBT, 0)), S.Context.getFloatTypeSemantics(QualType(SourceBT, 0)))) { if (S.getLangOpts().SYCLIsDevice) - S.SYCLDiagIfDeviceCode(CC, diag::warn_imp_float_size_conversion); + S.SYCL().DiagIfDeviceCode(CC, + diag::warn_imp_float_size_conversion); else DiagnoseImpCast(S, E, T, CC, diag::warn_imp_float_size_conversion); @@ -16588,7 +16589,7 @@ static void CheckImplicitConversion(Sema &S, Expr *E, QualType T, // warning. if (S.Diags.isIgnored(diag::warn_impcast_float_precision, CC)) { if (S.getLangOpts().SYCLIsDevice) - S.SYCLDiagIfDeviceCode(CC, diag::warn_imp_float_size_conversion); + S.SYCL().DiagIfDeviceCode(CC, diag::warn_imp_float_size_conversion); else DiagnoseImpCast(S, E, T, CC, diag::warn_imp_float_size_conversion); } diff --git a/clang/lib/Sema/SemaDecl.cpp b/clang/lib/Sema/SemaDecl.cpp index a5b0ebff72482..2ceefc12b56c3 100644 --- a/clang/lib/Sema/SemaDecl.cpp +++ b/clang/lib/Sema/SemaDecl.cpp @@ -46,6 +46,7 @@ #include "clang/Sema/Scope.h" #include "clang/Sema/ScopeInfo.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "clang/Sema/Template.h" #include "llvm/ADT/SmallString.h" #include "llvm/ADT/StringExtras.h" @@ -7972,13 +7973,13 @@ NamedDecl *Sema::ActOnVariableDeclarator( if (getLangOpts().SYCLIsDevice) { // device_global array is not allowed. if (const ArrayType *AT = getASTContext().getAsArrayType(NewVD->getType())) - if (isTypeDecoratedWithDeclAttribute( + if (SYCL().isTypeDecoratedWithDeclAttribute( AT->getElementType())) Diag(NewVD->getLocation(), diag::err_sycl_device_global_array); // Global variables with types decorated with device_global attribute must // be static if they are declared in SYCL device code. - if (isTypeDecoratedWithDeclAttribute( + if (SYCL().isTypeDecoratedWithDeclAttribute( NewVD->getType())) { if (SCSpec == DeclSpec::SCS_static) { const DeclContext *DC = NewVD->getDeclContext(); @@ -8000,10 +8001,10 @@ NamedDecl *Sema::ActOnVariableDeclarator( // constexpr unless their types are decorated with global_variable_allowed // attribute. if (SCSpec == DeclSpec::SCS_static && !R.isConstant(Context) && - !isTypeDecoratedWithDeclAttribute( + !SYCL().isTypeDecoratedWithDeclAttribute( NewVD->getType())) - SYCLDiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) - << Sema::KernelNonConstStaticDataVariable; + SYCL().DiagIfDeviceCode(D.getIdentifierLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelNonConstStaticDataVariable; } switch (D.getDeclSpec().getConstexprSpecifier()) { @@ -8168,8 +8169,8 @@ NamedDecl *Sema::ActOnVariableDeclarator( if (!Context.getTargetInfo().isValidGCCRegisterName(Label) && DeclAttrsMatchCUDAMode(getLangOpts(), getCurFunctionDecl())) { if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(E->getExprLoc(), - diag::err_asm_unknown_register_name) + SYCL().DiagIfDeviceCode(E->getExprLoc(), + diag::err_asm_unknown_register_name) << Label; else Diag(E->getExprLoc(), diag::err_asm_unknown_register_name) << Label; @@ -8362,7 +8363,7 @@ NamedDecl *Sema::ActOnVariableDeclarator( if (IsMemberSpecialization && !NewVD->isInvalidDecl()) CompleteMemberSpecialization(NewVD, Previous); - addSyclVarDecl(NewVD); + SYCL().addSyclVarDecl(NewVD); emitReadOnlyPlacementAttrWarning(*this, NewVD); return NewVD; @@ -13853,7 +13854,7 @@ void Sema::AddInitializerToDecl(Decl *RealDecl, Expr *Init, bool DirectInit) { } // In the SYCL explicit SIMD extension non constant "private globals" can't // be explicitly initialized in the declaration. - if (isSYCLEsimdPrivateGlobal(VDecl)) { + if (SYCL().isSYCLEsimdPrivateGlobal(VDecl)) { Diag(VDecl->getLocation(), diag::err_esimd_glob_cant_init); VDecl->setInvalidDecl(); return; @@ -14510,7 +14511,7 @@ void Sema::ActOnUninitializedDecl(Decl *RealDecl) { return; // In SYCL explicit SIMD extension "private global" variables can't be // initialized even implicitly, so don't synthesize an implicit initializer. - if (isSYCLEsimdPrivateGlobal(Var)) + if (SYCL().isSYCLEsimdPrivateGlobal(Var)) return; // C++03 [dcl.init]p9: @@ -14652,7 +14653,7 @@ void Sema::CheckCompleteVariableDeclaration(VarDecl *var) { } if (getLangOpts().SYCLIsDevice) - checkSYCLDeviceVarDecl(var); + SYCL().checkSYCLDeviceVarDecl(var); // In Objective-C, don't allow jumps past the implicit initialization of a // local retaining variable. diff --git a/clang/lib/Sema/SemaDeclAttr.cpp b/clang/lib/Sema/SemaDeclAttr.cpp index 94bbca32063cc..fb905a669c91f 100644 --- a/clang/lib/Sema/SemaDeclAttr.cpp +++ b/clang/lib/Sema/SemaDeclAttr.cpp @@ -40,6 +40,7 @@ #include "clang/Sema/Scope.h" #include "clang/Sema/ScopeInfo.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "llvm/ADT/STLExtras.h" #include "llvm/ADT/StringExtras.h" #include "llvm/IR/Assumptions.h" @@ -6587,7 +6588,7 @@ static void handleSYCLDeviceAttr(Sema &S, Decl *D, const ParsedAttr &AL) { // Diagnose only for non-dependent types since dependent type don't have // attributes applied on them ATM. if (!VarType->isDependentType() && - !S.isTypeDecoratedWithDeclAttribute( + !S.SYCL().isTypeDecoratedWithDeclAttribute( VD->getType())) { S.Diag(AL.getLoc(), diag::err_sycl_attribute_not_device_global) << AL; return; @@ -7475,7 +7476,7 @@ static bool CheckValidFPGAMemoryAttributesVar(Sema &S, Decl *D) { if (!(isa(D) || (VD->getKind() != Decl::ImplicitParam && VD->getKind() != Decl::NonTypeTemplateParm && - (S.isTypeDecoratedWithDeclAttribute( + (S.SYCL().isTypeDecoratedWithDeclAttribute( VD->getType()) || VD->getType().isConstQualified() || VD->getType().getAddressSpace() == LangAS::opencl_constant || @@ -8170,7 +8171,7 @@ void Sema::AddSYCLIntelPrivateCopiesAttr(Decl *D, const AttributeCommonInfo &CI, VD->getKind() != Decl::NonTypeTemplateParm && VD->getKind() != Decl::ParmVar && (VD->hasLocalStorage() || - isTypeDecoratedWithDeclAttribute( + SYCL().isTypeDecoratedWithDeclAttribute( VD->getType())))))) { Diag(CI.getLoc(), diag::err_fpga_attribute_invalid_decl) << CI; return; diff --git a/clang/lib/Sema/SemaDeclCXX.cpp b/clang/lib/Sema/SemaDeclCXX.cpp index 9c51bd128eaff..8e43b02cfd790 100644 --- a/clang/lib/Sema/SemaDeclCXX.cpp +++ b/clang/lib/Sema/SemaDeclCXX.cpp @@ -43,6 +43,7 @@ #include "clang/Sema/Scope.h" #include "clang/Sema/ScopeInfo.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "clang/Sema/Template.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/STLExtras.h" @@ -3776,7 +3777,7 @@ Sema::ActOnCXXMemberDeclarator(Scope *S, AccessSpecifier AS, Declarator &D, // attribute is accessed. if (getLangOpts().SYCLIsDevice) { if (auto Value = dyn_cast(Member)) { - if (isTypeDecoratedWithDeclAttribute( + if (SYCL().isTypeDecoratedWithDeclAttribute( Value->getType())) { if (Value->getAccess() == AS_private || Value->getAccess() == AS_protected) { diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 92e1b1d740a3c..5f48833c1f6f9 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -233,28 +233,28 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, bool IsRuntimeEvaluated = ExprEvalContexts.empty() || (!isUnevaluatedContext() && !isConstantEvaluatedContext()); - bool IsEsimdPrivateGlobal = isSYCLEsimdPrivateGlobal(VD); + bool IsEsimdPrivateGlobal = SYCL().isSYCLEsimdPrivateGlobal(VD); // Non-const statics are not allowed in SYCL except for ESIMD or with the // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->getStorageClass() == SC_Static && !VD->hasAttr() && - !isTypeDecoratedWithDeclAttribute( - VD->getType())) - SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) - << Sema::KernelNonConstStaticDataVariable; + !SemaSYCL::isTypeDecoratedWithDeclAttribute< + SYCLGlobalVariableAllowedAttr>(VD->getType())) + SYCL().DiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) + << SemaSYCL::KernelNonConstStaticDataVariable; // Non-const globals are not allowed in SYCL except for ESIMD or with the // SYCLGlobalVar or SYCLGlobalVariableAllowed attribute. else if (IsRuntimeEvaluated && !IsEsimdPrivateGlobal && !IsConst && VD->hasGlobalStorage() && !VD->hasAttr() && - !isTypeDecoratedWithDeclAttribute( - VD->getType())) - SYCLDiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) - << Sema::KernelGlobalVariable; + !SemaSYCL::isTypeDecoratedWithDeclAttribute< + SYCLGlobalVariableAllowedAttr>(VD->getType())) + SYCL().DiagIfDeviceCode(*Locs.begin(), diag::err_sycl_restrict) + << SemaSYCL::KernelGlobalVariable; // ESIMD globals cannot be used in a SYCL context. else if (IsRuntimeEvaluated && IsEsimdPrivateGlobal && VD->hasGlobalStorage()) - SYCLDiagIfDeviceCode(*Locs.begin(), + SYCL().DiagIfDeviceCode(*Locs.begin(), diag::err_esimd_global_in_sycl_context, Sema::DeviceDiagnosticReason::Sycl); } else if (auto *FDecl = dyn_cast(D)) { @@ -288,7 +288,7 @@ bool Sema::DiagnoseUseOfDecl(NamedDecl *D, ArrayRef Locs, !Id->getName().starts_with("__devicelib_ConvertFToBF16INTEL") && !Id->getName().starts_with("__assert_fail") && !isMsvcMathFn(Id->getName())) { - SYCLDiagIfDeviceCode( + SYCL().DiagIfDeviceCode( *Locs.begin(), diag::err_sycl_device_function_is_called_from_esimd, Sema::DeviceDiagnosticReason::Esimd); } @@ -3510,11 +3510,10 @@ static bool ShouldLookupResultBeMultiVersionOverload(const LookupResult &R) { ExprResult Sema::BuildDeclarationNameExpr(const CXXScopeSpec &SS, LookupResult &R, bool NeedsADL, - bool AcceptInvalidDecl, - bool NeedUnresolved) { + bool AcceptInvalidDecl) { // If this is a single, fully-resolved result and we don't need ADL, // just build an ordinary singleton decl ref. - if (!NeedUnresolved && !NeedsADL && R.isSingleResult() && + if (!NeedsADL && R.isSingleResult() && !R.getAsSingle() && !ShouldLookupResultBeMultiVersionOverload(R)) return BuildDeclarationNameExpr(SS, R.getLookupNameInfo(), R.getFoundDecl(), @@ -3863,65 +3862,6 @@ ExprResult Sema::BuildPredefinedExpr(SourceLocation Loc, SL); } -ExprResult Sema::BuildSYCLUniqueStableNameExpr(SourceLocation OpLoc, - SourceLocation LParen, - SourceLocation RParen, - TypeSourceInfo *TSI) { - return SYCLUniqueStableNameExpr::Create(Context, OpLoc, LParen, RParen, TSI); -} - -ExprResult Sema::ActOnSYCLUniqueStableNameExpr(SourceLocation OpLoc, - SourceLocation LParen, - SourceLocation RParen, - ParsedType ParsedTy) { - TypeSourceInfo *TSI = nullptr; - QualType Ty = GetTypeFromParser(ParsedTy, &TSI); - - if (Ty.isNull()) - return ExprError(); - if (!TSI) - TSI = Context.getTrivialTypeSourceInfo(Ty, LParen); - - return BuildSYCLUniqueStableNameExpr(OpLoc, LParen, RParen, TSI); -} - -ExprResult Sema::BuildSYCLUniqueStableIdExpr(SourceLocation OpLoc, - SourceLocation LParen, - SourceLocation RParen, Expr *E) { - if (!E->isInstantiationDependent()) { - // Special handling to get us better error messages for a member variable. - if (auto *ME = dyn_cast(E->IgnoreUnlessSpelledInSource())) { - if (isa(ME->getMemberDecl())) - Diag(E->getExprLoc(), diag::err_unique_stable_id_global_storage); - else - Diag(E->getExprLoc(), diag::err_unique_stable_id_expected_var); - return ExprError(); - } - - auto *DRE = dyn_cast(E->IgnoreUnlessSpelledInSource()); - - if (!DRE || !isa_and_nonnull(DRE->getDecl())) { - Diag(E->getExprLoc(), diag::err_unique_stable_id_expected_var); - return ExprError(); - } - - auto *Var = cast(DRE->getDecl()); - - if (!Var->hasGlobalStorage()) { - Diag(E->getExprLoc(), diag::err_unique_stable_id_global_storage); - return ExprError(); - } - } - - return SYCLUniqueStableIdExpr::Create(Context, OpLoc, LParen, RParen, E); -} - -ExprResult Sema::ActOnSYCLUniqueStableIdExpr(SourceLocation OpLoc, - SourceLocation LParen, - SourceLocation RParen, Expr *E) { - return BuildSYCLUniqueStableIdExpr(OpLoc, LParen, RParen, E); -} - ExprResult Sema::ActOnPredefinedExpr(SourceLocation Loc, tok::TokenKind Kind) { return BuildPredefinedExpr(Loc, getPredefinedExprKind(Kind)); } diff --git a/clang/lib/Sema/SemaExprCXX.cpp b/clang/lib/Sema/SemaExprCXX.cpp index 7c52e4ab70162..433211ef6790c 100644 --- a/clang/lib/Sema/SemaExprCXX.cpp +++ b/clang/lib/Sema/SemaExprCXX.cpp @@ -889,8 +889,8 @@ ExprResult Sema::BuildCXXThrow(SourceLocation OpLoc, Expr *Ex, // Exceptions aren't allowed in SYCL device code. if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(OpLoc, diag::err_sycl_restrict) - << Sema::KernelUseExceptions; + SYCL().DiagIfDeviceCode(OpLoc, diag::err_sycl_restrict) + << SemaSYCL::KernelUseExceptions; if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope()) Diag(OpLoc, diag::err_omp_simd_region_cannot_use_stmt) << "throw"; @@ -1420,42 +1420,26 @@ bool Sema::CheckCXXThisCapture(SourceLocation Loc, const bool Explicit, } ExprResult Sema::ActOnCXXThis(SourceLocation Loc) { - // C++20 [expr.prim.this]p1: - // The keyword this names a pointer to the object for which an - // implicit object member function is invoked or a non-static - // data member's initializer is evaluated. + /// C++ 9.3.2: In the body of a non-static member function, the keyword this + /// is a non-lvalue expression whose value is the address of the object for + /// which the function is called. QualType ThisTy = getCurrentThisType(); - if (CheckCXXThisType(Loc, ThisTy)) - return ExprError(); + if (ThisTy.isNull()) { + DeclContext *DC = getFunctionLevelDeclContext(); - return BuildCXXThisExpr(Loc, ThisTy, /*IsImplicit=*/false); -} + if (const auto *Method = dyn_cast(DC); + Method && Method->isExplicitObjectMemberFunction()) { + return Diag(Loc, diag::err_invalid_this_use) << 1; + } -bool Sema::CheckCXXThisType(SourceLocation Loc, QualType Type) { - if (!Type.isNull()) - return false; + if (isLambdaCallWithExplicitObjectParameter(CurContext)) + return Diag(Loc, diag::err_invalid_this_use) << 1; - // C++20 [expr.prim.this]p3: - // If a declaration declares a member function or member function template - // of a class X, the expression this is a prvalue of type - // "pointer to cv-qualifier-seq X" wherever X is the current class between - // the optional cv-qualifier-seq and the end of the function-definition, - // member-declarator, or declarator. It shall not appear within the - // declaration of either a static member function or an explicit object - // member function of the current class (although its type and value - // category are defined within such member functions as they are within - // an implicit object member function). - DeclContext *DC = getFunctionLevelDeclContext(); - if (const auto *Method = dyn_cast(DC); - Method && Method->isExplicitObjectMemberFunction()) { - Diag(Loc, diag::err_invalid_this_use) << 1; - } else if (isLambdaCallWithExplicitObjectParameter(CurContext)) { - Diag(Loc, diag::err_invalid_this_use) << 1; - } else { - Diag(Loc, diag::err_invalid_this_use) << 0; + return Diag(Loc, diag::err_invalid_this_use) << 0; } - return true; + + return BuildCXXThisExpr(Loc, ThisTy, /*IsImplicit=*/false); } Expr *Sema::BuildCXXThisExpr(SourceLocation Loc, QualType Type, @@ -2535,8 +2519,8 @@ ExprResult Sema::BuildCXXNew(SourceRange Range, bool UseGlobal, MarkFunctionReferenced(StartLoc, OperatorNew); if (getLangOpts().SYCLIsDevice && OperatorNew->isReplaceableGlobalAllocationFunction()) - SYCLDiagIfDeviceCode(StartLoc, diag::err_sycl_restrict) - << KernelAllocateStorage; + SYCL().DiagIfDeviceCode(StartLoc, diag::err_sycl_restrict) + << SemaSYCL::KernelAllocateStorage; } if (OperatorDelete) { if (DiagnoseUseOfDecl(OperatorDelete, StartLoc)) diff --git a/clang/lib/Sema/SemaExprMember.cpp b/clang/lib/Sema/SemaExprMember.cpp index 8cd2288d279cc..32998ae60eafe 100644 --- a/clang/lib/Sema/SemaExprMember.cpp +++ b/clang/lib/Sema/SemaExprMember.cpp @@ -61,10 +61,6 @@ enum IMAKind { /// The reference is a contextually-permitted abstract member reference. IMA_Abstract, - /// Whether the context is static is dependent on the enclosing template (i.e. - /// in a dependent class scope explicit specialization). - IMA_Dependent, - /// The reference may be to an unresolved using declaration and the /// context is not an instance method. IMA_Unresolved_StaticOrExplicitContext, @@ -95,18 +91,10 @@ static IMAKind ClassifyImplicitMemberAccess(Sema &SemaRef, DeclContext *DC = SemaRef.getFunctionLevelDeclContext(); - bool couldInstantiateToStatic = false; - bool isStaticOrExplicitContext = SemaRef.CXXThisTypeOverride.isNull(); - - if (auto *MD = dyn_cast(DC)) { - if (MD->isImplicitObjectMemberFunction()) { - isStaticOrExplicitContext = false; - // A dependent class scope function template explicit specialization - // that is neither declared 'static' nor with an explicit object - // parameter could instantiate to a static or non-static member function. - couldInstantiateToStatic = MD->getDependentSpecializationInfo(); - } - } + bool isStaticOrExplicitContext = + SemaRef.CXXThisTypeOverride.isNull() && + (!isa(DC) || cast(DC)->isStatic() || + cast(DC)->isExplicitObjectMemberFunction()); if (R.isUnresolvableResult()) return isStaticOrExplicitContext ? IMA_Unresolved_StaticOrExplicitContext @@ -135,9 +123,6 @@ static IMAKind ClassifyImplicitMemberAccess(Sema &SemaRef, if (Classes.empty()) return IMA_Static; - if (couldInstantiateToStatic) - return IMA_Dependent; - // C++11 [expr.prim.general]p12: // An id-expression that denotes a non-static data member or non-static // member function of a class can only be used: @@ -283,30 +268,27 @@ ExprResult Sema::BuildPossibleImplicitMemberExpr( const CXXScopeSpec &SS, SourceLocation TemplateKWLoc, LookupResult &R, const TemplateArgumentListInfo *TemplateArgs, const Scope *S, UnresolvedLookupExpr *AsULE) { - switch (IMAKind Classification = ClassifyImplicitMemberAccess(*this, R)) { + switch (ClassifyImplicitMemberAccess(*this, R)) { case IMA_Instance: + return BuildImplicitMemberExpr(SS, TemplateKWLoc, R, TemplateArgs, true, S); + case IMA_Mixed: case IMA_Mixed_Unrelated: case IMA_Unresolved: - return BuildImplicitMemberExpr( - SS, TemplateKWLoc, R, TemplateArgs, - /*IsKnownInstance=*/Classification == IMA_Instance, S); + return BuildImplicitMemberExpr(SS, TemplateKWLoc, R, TemplateArgs, false, + S); + case IMA_Field_Uneval_Context: Diag(R.getNameLoc(), diag::warn_cxx98_compat_non_static_member_use) << R.getLookupNameInfo().getName(); [[fallthrough]]; case IMA_Static: case IMA_Abstract: - case IMA_Dependent: case IMA_Mixed_StaticOrExplicitContext: case IMA_Unresolved_StaticOrExplicitContext: if (TemplateArgs || TemplateKWLoc.isValid()) - return BuildTemplateIdExpr(SS, TemplateKWLoc, R, /*RequiresADL=*/false, - TemplateArgs); - return AsULE ? AsULE - : BuildDeclarationNameExpr( - SS, R, /*NeedsADL=*/false, /*AcceptInvalidDecl=*/false, - /*NeedUnresolved=*/Classification == IMA_Dependent); + return BuildTemplateIdExpr(SS, TemplateKWLoc, R, false, TemplateArgs); + return AsULE ? AsULE : BuildDeclarationNameExpr(SS, R, false); case IMA_Error_StaticOrExplicitContext: case IMA_Error_Unrelated: diff --git a/clang/lib/Sema/SemaOverload.cpp b/clang/lib/Sema/SemaOverload.cpp index a1d16cc2c633b..ec941178258a0 100644 --- a/clang/lib/Sema/SemaOverload.cpp +++ b/clang/lib/Sema/SemaOverload.cpp @@ -32,6 +32,7 @@ #include "clang/Sema/Lookup.h" #include "clang/Sema/Overload.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "clang/Sema/Template.h" #include "clang/Sema/TemplateDeduction.h" #include "llvm/ADT/DenseSet.h" @@ -11110,9 +11111,9 @@ static bool checkAddressOfFunctionIsAvailable(Sema &S, const FunctionDecl *FD, if (Complain && S.getLangOpts().SYCLIsDevice && S.getLangOpts().SYCLAllowFuncPtr) { if (!FD->hasAttr()) { - S.SYCLDiagIfDeviceCode(Loc, - diag::err_sycl_taking_address_of_wrong_function, - Sema::DeviceDiagnosticReason::Sycl); + S.SYCL().DiagIfDeviceCode(Loc, + diag::err_sycl_taking_address_of_wrong_function, + Sema::DeviceDiagnosticReason::Sycl); } } diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index b1fb7f25066ab..80932bb56e78b 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -8,6 +8,7 @@ // This implements Semantic Analysis for SYCL constructs. //===----------------------------------------------------------------------===// +#include "clang/Sema/SemaSYCL.h" #include "TreeTransform.h" #include "clang/AST/AST.h" #include "clang/AST/Mangle.h" @@ -69,7 +70,7 @@ static constexpr llvm::StringLiteral LibstdcxxFailedAssertion = "__failed_assertion"; constexpr unsigned MaxKernelArgsSize = 2048; -bool Sema::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) { +bool SemaSYCL::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) { const auto *RD = Ty->getAsCXXRecordDecl(); if (!RD) return false; @@ -87,8 +88,8 @@ bool Sema::isSyclType(QualType Ty, SYCLTypeAttr::SYCLType TypeName) { } static bool isSyclAccessorType(QualType Ty) { - return Sema::isSyclType(Ty, SYCLTypeAttr::accessor) || - Sema::isSyclType(Ty, SYCLTypeAttr::local_accessor); + return SemaSYCL::isSyclType(Ty, SYCLTypeAttr::accessor) || + SemaSYCL::isSyclType(Ty, SYCLTypeAttr::local_accessor); } // FIXME: Accessor property lists should be modified to use compile-time @@ -105,25 +106,25 @@ static bool isAccessorPropertyType(QualType Ty, return false; } -static bool isSyclSpecialType(QualType Ty, Sema &S) { +static bool isSyclSpecialType(QualType Ty, SemaSYCL &S) { return S.isTypeDecoratedWithDeclAttribute(Ty); } -ExprResult Sema::ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT) { +ExprResult SemaSYCL::ActOnSYCLBuiltinNumFieldsExpr(ParsedType PT) { TypeSourceInfo *TInfo = nullptr; - QualType QT = GetTypeFromParser(PT, &TInfo); + QualType QT = Sema::GetTypeFromParser(PT, &TInfo); assert(TInfo && "couldn't get type info from a type from the parser?"); SourceLocation TypeLoc = TInfo->getTypeLoc().getBeginLoc(); return BuildSYCLBuiltinNumFieldsExpr(TypeLoc, QT); } -ExprResult Sema::BuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, - QualType SourceTy) { +ExprResult SemaSYCL::BuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, + QualType SourceTy) { if (!SourceTy->isDependentType()) { - if (RequireCompleteType(Loc, SourceTy, - diag::err_sycl_type_trait_requires_complete_type, - /*__builtin_num_fields*/ 0)) + if (SemaRef.RequireCompleteType( + Loc, SourceTy, diag::err_sycl_type_trait_requires_complete_type, + /*__builtin_num_fields*/ 0)) return ExprError(); if (!SourceTy->isRecordType()) { @@ -132,24 +133,25 @@ ExprResult Sema::BuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, return ExprError(); } } - return new (Context) - SYCLBuiltinNumFieldsExpr(Loc, SourceTy, Context.getSizeType()); + return new (getASTContext()) + SYCLBuiltinNumFieldsExpr(Loc, SourceTy, getASTContext().getSizeType()); } -ExprResult Sema::ActOnSYCLBuiltinFieldTypeExpr(ParsedType PT, Expr *Idx) { +ExprResult SemaSYCL::ActOnSYCLBuiltinFieldTypeExpr(ParsedType PT, Expr *Idx) { TypeSourceInfo *TInfo = nullptr; - QualType QT = GetTypeFromParser(PT, &TInfo); + QualType QT = Sema::GetTypeFromParser(PT, &TInfo); assert(TInfo && "couldn't get type info from a type from the parser?"); SourceLocation TypeLoc = TInfo->getTypeLoc().getBeginLoc(); return BuildSYCLBuiltinFieldTypeExpr(TypeLoc, QT, Idx); } -ExprResult Sema::BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, - QualType SourceTy, Expr *Idx) { +ExprResult SemaSYCL::BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, + QualType SourceTy, + Expr *Idx) { // If the expression appears in an evaluated context, we want to give an // error so that users don't attempt to use the value of this expression. - if (!isUnevaluatedContext()) { + if (!SemaRef.isUnevaluatedContext()) { Diag(Loc, diag::err_sycl_builtin_type_trait_evaluated) << /*__builtin_field_type*/ 0; return ExprError(); @@ -162,9 +164,9 @@ ExprResult Sema::BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, QualType FieldTy = SourceTy; ExprValueKind ValueKind = VK_PRValue; if (!SourceTy->isDependentType()) { - if (RequireCompleteType(Loc, SourceTy, - diag::err_sycl_type_trait_requires_complete_type, - /*__builtin_field_type*/ 1)) + if (SemaRef.RequireCompleteType( + Loc, SourceTy, diag::err_sycl_type_trait_requires_complete_type, + /*__builtin_field_type*/ 1)) return ExprError(); if (!SourceTy->isRecordType()) { @@ -174,7 +176,8 @@ ExprResult Sema::BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, } if (!Idx->isValueDependent()) { - std::optional IdxVal = Idx->getIntegerConstantExpr(Context); + std::optional IdxVal = + Idx->getIntegerConstantExpr(getASTContext()); if (IdxVal) { RecordDecl *RD = SourceTy->getAsRecordDecl(); assert(RD && "Record type but no record decl?"); @@ -209,25 +212,25 @@ ExprResult Sema::BuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, } } } - return new (Context) + return new (getASTContext()) SYCLBuiltinFieldTypeExpr(Loc, SourceTy, Idx, FieldTy, ValueKind); } -ExprResult Sema::ActOnSYCLBuiltinNumBasesExpr(ParsedType PT) { +ExprResult SemaSYCL::ActOnSYCLBuiltinNumBasesExpr(ParsedType PT) { TypeSourceInfo *TInfo = nullptr; - QualType QT = GetTypeFromParser(PT, &TInfo); + QualType QT = Sema::GetTypeFromParser(PT, &TInfo); assert(TInfo && "couldn't get type info from a type from the parser?"); SourceLocation TypeLoc = TInfo->getTypeLoc().getBeginLoc(); return BuildSYCLBuiltinNumBasesExpr(TypeLoc, QT); } -ExprResult Sema::BuildSYCLBuiltinNumBasesExpr(SourceLocation Loc, +ExprResult SemaSYCL::BuildSYCLBuiltinNumBasesExpr(SourceLocation Loc, QualType SourceTy) { if (!SourceTy->isDependentType()) { - if (RequireCompleteType(Loc, SourceTy, - diag::err_sycl_type_trait_requires_complete_type, - /*__builtin_num_bases*/ 2)) + if (SemaRef.RequireCompleteType( + Loc, SourceTy, diag::err_sycl_type_trait_requires_complete_type, + /*__builtin_num_bases*/ 2)) return ExprError(); if (!SourceTy->isRecordType()) { @@ -236,24 +239,25 @@ ExprResult Sema::BuildSYCLBuiltinNumBasesExpr(SourceLocation Loc, return ExprError(); } } - return new (Context) - SYCLBuiltinNumBasesExpr(Loc, SourceTy, Context.getSizeType()); + return new (getASTContext()) + SYCLBuiltinNumBasesExpr(Loc, SourceTy, getASTContext().getSizeType()); } -ExprResult Sema::ActOnSYCLBuiltinBaseTypeExpr(ParsedType PT, Expr *Idx) { +ExprResult SemaSYCL::ActOnSYCLBuiltinBaseTypeExpr(ParsedType PT, Expr *Idx) { TypeSourceInfo *TInfo = nullptr; - QualType QT = GetTypeFromParser(PT, &TInfo); + QualType QT = SemaRef.GetTypeFromParser(PT, &TInfo); assert(TInfo && "couldn't get type info from a type from the parser?"); SourceLocation TypeLoc = TInfo->getTypeLoc().getBeginLoc(); return BuildSYCLBuiltinBaseTypeExpr(TypeLoc, QT, Idx); } -ExprResult Sema::BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, - QualType SourceTy, Expr *Idx) { +ExprResult SemaSYCL::BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, + QualType SourceTy, + Expr *Idx) { // If the expression appears in an evaluated context, we want to give an // error so that users don't attempt to use the value of this expression. - if (!isUnevaluatedContext()) { + if (!SemaRef.isUnevaluatedContext()) { Diag(Loc, diag::err_sycl_builtin_type_trait_evaluated) << /*__builtin_base_type*/ 1; return ExprError(); @@ -265,9 +269,9 @@ ExprResult Sema::BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, // later to the real type. QualType BaseTy = SourceTy; if (!SourceTy->isDependentType()) { - if (RequireCompleteType(Loc, SourceTy, - diag::err_sycl_type_trait_requires_complete_type, - /*__builtin_base_type*/ 3)) + if (SemaRef.RequireCompleteType( + Loc, SourceTy, diag::err_sycl_type_trait_requires_complete_type, + /*__builtin_base_type*/ 3)) return ExprError(); if (!SourceTy->isRecordType()) { @@ -277,7 +281,8 @@ ExprResult Sema::BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, } if (!Idx->isValueDependent()) { - std::optional IdxVal = Idx->getIntegerConstantExpr(Context); + std::optional IdxVal = + Idx->getIntegerConstantExpr(getASTContext()); if (IdxVal) { CXXRecordDecl *RD = SourceTy->getAsCXXRecordDecl(); assert(RD && "Record type but no record decl?"); @@ -303,7 +308,8 @@ ExprResult Sema::BuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, } } } - return new (Context) SYCLBuiltinBaseTypeExpr(Loc, SourceTy, Idx, BaseTy); + return new (getASTContext()) + SYCLBuiltinBaseTypeExpr(Loc, SourceTy, Idx, BaseTy); } // This information is from Section 4.13 of the SYCL spec @@ -366,7 +372,7 @@ static bool IsSyclMathFunc(unsigned BuiltinID) { return true; } -bool Sema::isDeclAllowedInSYCLDeviceCode(const Decl *D) { +bool SemaSYCL::isDeclAllowedInSYCLDeviceCode(const Decl *D) { if (const FunctionDecl *FD = dyn_cast(D)) { const IdentifierInfo *II = FD->getIdentifier(); @@ -378,7 +384,7 @@ bool Sema::isDeclAllowedInSYCLDeviceCode(const Decl *D) { return true; // Allow to use `::printf` only for CUDA. - if (Context.getTargetInfo().getTriple().isNVPTX()) { + if (getASTContext().getTargetInfo().getTriple().isNVPTX()) { if (FD->getBuiltinID() == Builtin::BIprintf) return true; } @@ -392,13 +398,16 @@ bool Sema::isDeclAllowedInSYCLDeviceCode(const Decl *D) { return false; } -static bool isZeroSizedArray(Sema &SemaRef, QualType Ty) { - if (const auto *CAT = SemaRef.getASTContext().getAsConstantArrayType(Ty)) +SemaSYCL::SemaSYCL(Sema &S) + : SemaBase(S), SyclIntHeader(nullptr), SyclIntFooter(nullptr) {} + +static bool isZeroSizedArray(SemaSYCL &S, QualType Ty) { + if (const auto *CAT = S.getASTContext().getAsConstantArrayType(Ty)) return CAT->isZeroSize(); return false; } -static void checkSYCLType(Sema &S, QualType Ty, SourceRange Loc, +static void checkSYCLType(SemaSYCL &S, QualType Ty, SourceRange Loc, llvm::DenseSet Visited, SourceRange UsedAtLoc = SourceRange()) { // Not all variable types are supported inside SYCL kernels, @@ -417,14 +426,15 @@ static void checkSYCLType(Sema &S, QualType Ty, SourceRange Loc, // zero length arrays if (isZeroSizedArray(S, Ty)) { - S.SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_typecheck_zero_array_size) + S.DiagIfDeviceCode(Loc.getBegin(), + diag::err_typecheck_zero_array_size) << 1; Emitting = true; } // variable length arrays if (Ty->isVariableArrayType()) { - S.SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_vla_unsupported) << 0; + S.DiagIfDeviceCode(Loc.getBegin(), diag::err_vla_unsupported) << 0; Emitting = true; } @@ -438,14 +448,14 @@ static void checkSYCLType(Sema &S, QualType Ty, SourceRange Loc, Ty->isSpecificBuiltinType(BuiltinType::LongDouble) || Ty->isSpecificBuiltinType(BuiltinType::BFloat16) || (Ty->isSpecificBuiltinType(BuiltinType::Float128) && - !S.Context.getTargetInfo().hasFloat128Type())) { - S.SYCLDiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) + !S.getASTContext().getTargetInfo().hasFloat128Type())) { + S.DiagIfDeviceCode(Loc.getBegin(), diag::err_type_unsupported) << Ty.getUnqualifiedType().getCanonicalType(); Emitting = true; } if (Emitting && UsedAtLoc.isValid()) - S.SYCLDiagIfDeviceCode(UsedAtLoc.getBegin(), diag::note_used_here); + S.DiagIfDeviceCode(UsedAtLoc.getBegin(), diag::note_used_here); //--- now recurse --- // Pointers complicate recursion. Add this type to Visited. @@ -466,7 +476,7 @@ static void checkSYCLType(Sema &S, QualType Ty, SourceRange Loc, } } -void Sema::checkSYCLDeviceVarDecl(VarDecl *Var) { +void SemaSYCL::checkSYCLDeviceVarDecl(VarDecl *Var) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); QualType Ty = Var->getType(); @@ -506,7 +516,7 @@ static bool isSYCLUndefinedAllowed(const FunctionDecl *Callee, // Helper function to report conflicting function attributes. // F - the function, A1 - function attribute, A2 - the attribute it conflicts // with. -static void reportConflictingAttrs(Sema &S, FunctionDecl *F, const Attr *A1, +static void reportConflictingAttrs(SemaSYCL &S, FunctionDecl *F, const Attr *A1, const Attr *A2) { S.Diag(F->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(A1->getLocation(), diag::note_conflicting_attribute); @@ -520,7 +530,7 @@ static int64_t getIntExprValue(const Expr *E, ASTContext &Ctx) { } // Collect function attributes related to SYCL. -static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, +static void collectSYCLAttributes(SemaSYCL &S, FunctionDecl *FD, llvm::SmallVectorImpl &Attrs, bool DirectlyCalled) { if (!FD->hasAttrs()) @@ -557,14 +567,14 @@ static void collectSYCLAttributes(Sema &S, FunctionDecl *FD, } class DiagDeviceFunction : public RecursiveASTVisitor { - Sema &SemaRef; + SemaSYCL &SemaSYCLRef; const llvm::SmallPtrSetImpl &RecursiveFuncs; public: DiagDeviceFunction( - Sema &S, + SemaSYCL &S, const llvm::SmallPtrSetImpl &RecursiveFuncs) - : RecursiveASTVisitor(), SemaRef(S), RecursiveFuncs(RecursiveFuncs) {} + : RecursiveASTVisitor(), SemaSYCLRef(S), RecursiveFuncs(RecursiveFuncs) {} void CheckBody(Stmt *ToBeDiagnosed) { TraverseStmt(ToBeDiagnosed); } @@ -578,18 +588,18 @@ class DiagDeviceFunction : public RecursiveASTVisitor { // all functions used by kernel have already been parsed and have // definitions. if (RecursiveFuncs.count(Callee)) { - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) - << Sema::KernelCallRecursiveFunction; - SemaRef.Diag(Callee->getSourceRange().getBegin(), - diag::note_sycl_recursive_function_declared_here) - << Sema::KernelCallRecursiveFunction; + SemaSYCLRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelCallRecursiveFunction; + SemaSYCLRef.Diag(Callee->getSourceRange().getBegin(), + diag::note_sycl_recursive_function_declared_here) + << SemaSYCL::KernelCallRecursiveFunction; } if (const CXXMethodDecl *Method = dyn_cast(Callee)) if (Method->isVirtual() && - !SemaRef.getLangOpts().SYCLAllowVirtualFunctions) - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) - << Sema::KernelCallVirtualFunction; + !SemaSYCLRef.getLangOpts().SYCLAllowVirtualFunctions) + SemaSYCLRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelCallVirtualFunction; if (auto const *FD = dyn_cast(Callee)) { // FIXME: We need check all target specified attributes for error if @@ -598,41 +608,43 @@ class DiagDeviceFunction : public RecursiveASTVisitor { // currently. Erich is currently working on that in LLVM, once that is // committed we need to change this". if (FD->hasAttr()) { - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) - << Sema::KernelCallDllimportFunction; - SemaRef.Diag(FD->getLocation(), diag::note_callee_decl) << FD; + SemaSYCLRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelCallDllimportFunction; + SemaSYCLRef.Diag(FD->getLocation(), diag::note_callee_decl) << FD; } } // Specifically check if the math library function corresponding to this // builtin is supported for SYCL unsigned BuiltinID = Callee->getBuiltinID(); if (BuiltinID && !IsSyclMathFunc(BuiltinID)) { - StringRef Name = SemaRef.Context.BuiltinInfo.getName(BuiltinID); - SemaRef.Diag(e->getExprLoc(), diag::err_builtin_target_unsupported) + StringRef Name = SemaSYCLRef.getASTContext().BuiltinInfo.getName(BuiltinID); + SemaSYCLRef.Diag(e->getExprLoc(), diag::err_builtin_target_unsupported) << Name << "SYCL device"; } - } else if (!SemaRef.getLangOpts().SYCLAllowFuncPtr && + } else if (!SemaSYCLRef.getLangOpts().SYCLAllowFuncPtr && !e->isTypeDependent() && !isa(e->getCallee())) { bool MaybeConstantExpr = false; Expr *NonDirectCallee = e->getCallee(); if (!NonDirectCallee->isValueDependent()) MaybeConstantExpr = - NonDirectCallee->isCXX11ConstantExpr(SemaRef.getASTContext()); + NonDirectCallee->isCXX11ConstantExpr(SemaSYCLRef.getASTContext()); if (!MaybeConstantExpr) - SemaRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) - << Sema::KernelCallFunctionPointer; + SemaSYCLRef.Diag(e->getExprLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelCallFunctionPointer; } return true; } bool VisitCXXTypeidExpr(CXXTypeidExpr *E) { - SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << Sema::KernelRTTI; + SemaSYCLRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelRTTI; return true; } bool VisitCXXDynamicCastExpr(const CXXDynamicCastExpr *E) { - SemaRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) << Sema::KernelRTTI; + SemaSYCLRef.Diag(E->getExprLoc(), diag::err_sycl_restrict) + << SemaSYCL::KernelRTTI; return true; } @@ -667,7 +679,7 @@ class DiagDeviceFunction : public RecursiveASTVisitor { bool TraverseIfStmt(IfStmt *S) { if (std::optional ActiveStmt = - S->getNondiscardedCase(SemaRef.Context)) { + S->getNondiscardedCase(SemaSYCLRef.getASTContext())) { if (*ActiveStmt) return TraverseStmt(*ActiveStmt); return true; @@ -686,7 +698,7 @@ class DiagDeviceFunction : public RecursiveASTVisitor { class DeviceFunctionTracker { friend class SingleDeviceFunctionTracker; CallGraph CG; - Sema &SemaRef; + SemaSYCL &SemaSYCLRef; // The list of functions used on the device, kept so we can diagnose on them // later. llvm::SmallPtrSet DeviceFunctions; @@ -696,7 +708,7 @@ class DeviceFunctionTracker { for (CallGraphNode::CallRecord Record : CG.getRoot()->callees()) if (auto *FD = dyn_cast(Record.Callee->getDecl())) if (FD->hasBody() && FD->hasAttr()) - SemaRef.addSyclDeviceDecl(FD); + SemaSYCLRef.addSyclDeviceDecl(FD); } CallGraphNode *getNodeForKernel(FunctionDecl *Kernel) { @@ -712,14 +724,14 @@ class DeviceFunctionTracker { } public: - DeviceFunctionTracker(Sema &S) : SemaRef(S) { - CG.setSkipConstantExpressions(S.Context); + DeviceFunctionTracker(SemaSYCL &S) : SemaSYCLRef(S) { + CG.setSkipConstantExpressions(S.getASTContext()); CG.addToCallGraph(S.getASTContext().getTranslationUnitDecl()); CollectSyclExternalFuncs(); } ~DeviceFunctionTracker() { - DiagDeviceFunction Diagnoser{SemaRef, RecursiveFunctions}; + DiagDeviceFunction Diagnoser{SemaSYCLRef, RecursiveFunctions}; for (const FunctionDecl *FD : DeviceFunctions) if (const FunctionDecl *Def = FD->getDefinition()) Diagnoser.CheckBody(Def->getBody()); @@ -786,7 +798,7 @@ class SingleDeviceFunctionTracker { // cases later in finalizeSYCLDelayedAnalysis(). if (!CurrentDecl->isDefined() && !CurrentDecl->hasAttr() && !CurrentDecl->hasAttr()) - Parent.SemaRef.addFDToReachableFromSyclDevice(CurrentDecl, + Parent.SemaSYCLRef.addFDToReachableFromSyclDevice(CurrentDecl, CallStack.back()); // If this is a parallel_for_work_item that is declared in the @@ -799,7 +811,7 @@ class SingleDeviceFunctionTracker { isDeclaredInSYCLNamespace(CurrentDecl) && !CurrentDecl->hasAttr()) { CurrentDecl->addAttr( - SYCLScopeAttr::CreateImplicit(Parent.SemaRef.getASTContext(), + SYCLScopeAttr::CreateImplicit(Parent.SemaSYCLRef.getASTContext(), SYCLScopeAttr::Level::WorkItem)); } @@ -824,7 +836,7 @@ class SingleDeviceFunctionTracker { // Collect attributes for functions that aren't the root kernel. if (!CallStack.empty()) { bool DirectlyCalled = CallStack.size() == 1; - collectSYCLAttributes(Parent.SemaRef, CurrentDecl, CollectedAttributes, + collectSYCLAttributes(Parent.SemaSYCLRef, CurrentDecl, CollectedAttributes, DirectlyCalled); } @@ -882,7 +894,7 @@ class SingleDeviceFunctionTracker { // Always inline the KernelBody in the kernel entry point. For ESIMD // inlining is handled later down the pipeline. if (KernelBody && - Parent.SemaRef.getLangOpts().SYCLForceInlineKernelLambda && + Parent.SemaSYCLRef.getLangOpts().SYCLForceInlineKernelLambda && !KernelBody->hasAttr() && !KernelBody->hasAttr() && !KernelBody->hasAttr()) { @@ -948,7 +960,7 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor { // not a direct call - continue search return true; QualType Ty = Ctx.getRecordType(Call->getRecordDecl()); - if (!Sema::isSyclType(Ty, SYCLTypeAttr::group)) + if (!SemaSYCL::isSyclType(Ty, SYCLTypeAttr::group)) // not a member of sycl::group - continue search return true; auto Name = Callee->getName(); @@ -967,7 +979,7 @@ class MarkWIScopeFnVisitor : public RecursiveASTVisitor { }; static bool isSYCLPrivateMemoryVar(VarDecl *VD) { - return Sema::isSyclType(VD->getType(), SYCLTypeAttr::private_memory); + return SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::private_memory); } static void addScopeAttrToLocalVars(CXXMethodDecl &F) { @@ -1046,7 +1058,7 @@ static ParamDesc makeParamDesc(ASTContext &Ctx, StringRef Name, QualType Ty) { /// \return the target of given SYCL accessor type static target getAccessTarget(QualType FieldTy, const ClassTemplateSpecializationDecl *AccTy) { - if (Sema::isSyclType(FieldTy, SYCLTypeAttr::local_accessor)) + if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::local_accessor)) return local; return static_cast( @@ -1066,7 +1078,7 @@ static QualType calculateKernelNameType(ASTContext &Ctx, // Gets a name for the OpenCL kernel function, calculated from the first // template argument of the kernel caller function. static std::pair -constructKernelName(Sema &S, const FunctionDecl *KernelCallerFunc, +constructKernelName(SemaSYCL &S, const FunctionDecl *KernelCallerFunc, MangleContext &MC) { QualType KernelNameType = calculateKernelNameType(S.getASTContext(), KernelCallerFunc); @@ -1108,7 +1120,7 @@ static ParmVarDecl *getSyclKernelHandlerArg(FunctionDecl *KernelCallerFunc) { // Specialization constants in SYCL 2020 are not captured by lambda and // accessed through new optional lambda argument kernel_handler auto IsHandlerLambda = [](ParmVarDecl *PVD) { - return Sema::isSyclType(PVD->getType(), SYCLTypeAttr::kernel_handler); + return SemaSYCL::isSyclType(PVD->getType(), SYCLTypeAttr::kernel_handler); }; assert(llvm::count_if(KernelCallerFunc->parameters(), IsHandlerLambda) <= 1 && @@ -1149,7 +1161,7 @@ template <> struct bind_param { using type = FieldDecl *; }; template using bind_param_t = typename bind_param::type; class KernelObjVisitor { - Sema &SemaRef; + SemaSYCL &SemaSYCLRef; template void VisitUnionImpl(const CXXRecordDecl *Owner, ParentTy &Parent, @@ -1230,7 +1242,7 @@ class KernelObjVisitor { for (const auto &Base : Range) { QualType BaseTy = Base.getType(); // Handle accessor class as base - if (isSyclSpecialType(BaseTy, SemaRef)) + if (isSyclSpecialType(BaseTy, SemaSYCLRef)) (void)std::initializer_list{ (Handlers.handleSyclSpecialType(Owner, Base, BaseTy), 0)...}; else @@ -1289,7 +1301,7 @@ class KernelObjVisitor { return; const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(ArrayTy); + SemaSYCLRef.getASTContext().getAsConstantArrayType(ArrayTy); assert(CAT && "Should only be called on constant-size array."); QualType ET = CAT->getElementType(); uint64_t ElemCount = CAT->getSize().getZExtValue(); @@ -1308,7 +1320,7 @@ class KernelObjVisitor { template void visitField(const CXXRecordDecl *Owner, FieldDecl *Field, QualType FieldTy, HandlerTys &... Handlers) { - if (isSyclSpecialType(FieldTy, SemaRef)) + if (isSyclSpecialType(FieldTy, SemaSYCLRef)) KF_FOR_EACH(handleSyclSpecialType, Field, FieldTy); else if (FieldTy->isStructureOrClassType()) { if (KF_FOR_EACH(handleStructType, Field, FieldTy)) { @@ -1333,7 +1345,7 @@ class KernelObjVisitor { } public: - KernelObjVisitor(Sema &S) : SemaRef(S) {} + KernelObjVisitor(SemaSYCL &S) : SemaSYCLRef(S) {} template void VisitRecordBases(const CXXRecordDecl *KernelFunctor, @@ -1443,15 +1455,15 @@ class SyclKernelFieldHandlerBase { // universally required data). class SyclKernelFieldHandler : public SyclKernelFieldHandlerBase { protected: - Sema &SemaRef; - SyclKernelFieldHandler(Sema &S) : SemaRef(S) {} + SemaSYCL &SemaSYCLRef; + SyclKernelFieldHandler(SemaSYCL &S) : SemaSYCLRef(S) {} // Returns 'true' if the thing we're visiting (Based on the FD/QualType pair) // is an element of an array. FD will always be the array field. When // traversing the array field, Ty will be the type of the array field or the // type of array element (or some decomposed type from array). bool isArrayElement(const FieldDecl *FD, QualType Ty) const { - return !SemaRef.getASTContext().hasSameType(FD->getType(), Ty); + return !SemaSYCLRef.getASTContext().hasSameType(FD->getType(), Ty); } }; @@ -1611,7 +1623,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { bool checkNotCopyableToKernel(const FieldDecl *FD, QualType FieldTy) { if (FieldTy->isArrayType()) { if (const auto *CAT = - SemaRef.getASTContext().getAsConstantArrayType(FieldTy)) { + SemaSYCLRef.getASTContext().getAsConstantArrayType(FieldTy)) { QualType ET = CAT->getElementType(); return checkNotCopyableToKernel(FD, ET); } @@ -1625,24 +1637,24 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { bool checkPropertyListType(TemplateArgument PropList, SourceLocation Loc) { if (PropList.getKind() != TemplateArgument::ArgKind::Type) - return SemaRef.Diag( + return SemaSYCLRef.Diag( Loc, diag::err_sycl_invalid_accessor_property_template_param); QualType PropListTy = PropList.getAsType(); - if (!Sema::isSyclType(PropListTy, SYCLTypeAttr::accessor_property_list)) - return SemaRef.Diag( + if (!SemaSYCL::isSyclType(PropListTy, SYCLTypeAttr::accessor_property_list)) + return SemaSYCLRef.Diag( Loc, diag::err_sycl_invalid_accessor_property_template_param); const auto *AccPropListDecl = cast(PropListTy->getAsRecordDecl()); if (AccPropListDecl->getTemplateArgs().size() != 1) - return SemaRef.Diag(Loc, + return SemaSYCLRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number) << "accessor_property_list"; const auto TemplArg = AccPropListDecl->getTemplateArgs()[0]; if (TemplArg.getKind() != TemplateArgument::ArgKind::Pack) - return SemaRef.Diag( + return SemaSYCLRef.Diag( Loc, diag::err_sycl_invalid_accessor_property_list_template_param) << /*accessor_property_list*/ 0 << /*parameter pack*/ 0; @@ -1650,7 +1662,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { for (TemplateArgument::pack_iterator Prop = TemplArg.pack_begin(); Prop != TemplArg.pack_end(); ++Prop) { if (Prop->getKind() != TemplateArgument::ArgKind::Type) - return SemaRef.Diag( + return SemaSYCLRef.Diag( Loc, diag::err_sycl_invalid_accessor_property_list_template_param) << /*accessor_property_list pack argument*/ 1 << /*type*/ 1; @@ -1666,20 +1678,20 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { const auto *PropDecl = cast(PropTy->getAsRecordDecl()); if (PropDecl->getTemplateArgs().size() != 1) - return SemaRef.Diag(Loc, + return SemaSYCLRef.Diag(Loc, diag::err_sycl_invalid_property_list_param_number) << "buffer_location"; const auto BufferLoc = PropDecl->getTemplateArgs()[0]; if (BufferLoc.getKind() != TemplateArgument::ArgKind::Integral) - return SemaRef.Diag( + return SemaSYCLRef.Diag( Loc, diag::err_sycl_invalid_accessor_property_list_template_param) << /*buffer_location*/ 2 << /*non-negative integer*/ 2; int LocationID = static_cast(BufferLoc.getAsIntegral().getExtValue()); if (LocationID < 0) - return SemaRef.Diag( + return SemaSYCLRef.Diag( Loc, diag::err_sycl_invalid_accessor_property_list_template_param) << /*buffer_location*/ 2 << /*non-negative integer*/ 2; @@ -1688,15 +1700,15 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } bool checkSyclSpecialType(QualType Ty, SourceRange Loc) { - assert(isSyclSpecialType(Ty, SemaRef) && + assert(isSyclSpecialType(Ty, SemaSYCLRef) && "Should only be called on sycl special class types."); // Annotated pointers and annotated arguments must be captured // directly by the SYCL kernel. - if ((Sema::isSyclType(Ty, SYCLTypeAttr::annotated_ptr) || - Sema::isSyclType(Ty, SYCLTypeAttr::annotated_arg)) && + if ((SemaSYCL::isSyclType(Ty, SYCLTypeAttr::annotated_ptr) || + SemaSYCL::isSyclType(Ty, SYCLTypeAttr::annotated_arg)) && (StructFieldDepth > 0 || StructBaseDepth > 0)) - return SemaRef.Diag(Loc.getBegin(), + return SemaSYCLRef.Diag(Loc.getBegin(), diag::err_bad_kernel_param_data_members) << Ty << /*Struct*/ 1; @@ -1709,7 +1721,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { // Parameter packs are used by properties so they are always valid. if (TA.getKind() != TemplateArgument::Pack) { llvm::DenseSet Visited; - checkSYCLType(SemaRef, TA.getAsType(), Loc, Visited); + checkSYCLType(SemaSYCLRef, TA.getAsType(), Loc, Visited); } if (TAL.size() > 5) @@ -1719,7 +1731,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { } public: - SyclKernelFieldChecker(Sema &S) + SyclKernelFieldChecker(SemaSYCL &S) : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} static constexpr const bool VisitNthArrayElement = false; bool isValid() { return !IsInvalid; } @@ -1737,7 +1749,7 @@ class SyclKernelFieldChecker : public SyclKernelFieldHandler { if (RD->isLambda()) { for (const LambdaCapture &LC : RD->captures()) if (LC.capturesThis() && LC.isImplicit()) { - SemaRef.Diag(LC.getLocation(), diag::err_implicit_this_capture); + SemaSYCLRef.Diag(LC.getLocation(), diag::err_implicit_this_capture); IsInvalid = true; } } @@ -1809,7 +1821,7 @@ class SyclKernelUnionChecker : public SyclKernelFieldHandler { DiagnosticsEngine &Diag; public: - SyclKernelUnionChecker(Sema &S) + SyclKernelUnionChecker(SemaSYCL &S) : SyclKernelFieldHandler(S), Diag(S.getASTContext().getDiagnostics()) {} bool isValid() { return !IsInvalid; } static constexpr const bool VisitUnionBody = true; @@ -1856,7 +1868,7 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { static constexpr const bool VisitUnionBody = false; static constexpr const bool VisitNthArrayElement = false; - SyclKernelDecompMarker(Sema &S) : SyclKernelFieldHandler(S) { + SyclKernelDecompMarker(SemaSYCL &S) : SyclKernelFieldHandler(S) { // In order to prevent checking this over and over, just add a dummy-base // entry. CollectionStack.push_back(true); @@ -1894,14 +1906,14 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { if (CollectionStack.pop_back_val()) { if (!RD->hasAttr()) RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( - SemaRef.getASTContext())); + SemaSYCLRef.getASTContext())); CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { PointerStack.back() = true; if (!RD->hasAttr()) RD->addAttr( - SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + SYCLGenerateNewTypeAttr::CreateImplicit(SemaSYCLRef.getASTContext())); } return true; } @@ -1924,14 +1936,14 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { if (CollectionStack.pop_back_val()) { if (!RD->hasAttr()) RD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( - SemaRef.getASTContext())); + SemaSYCLRef.getASTContext())); CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { PointerStack.back() = true; if (!RD->hasAttr()) RD->addAttr( - SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + SYCLGenerateNewTypeAttr::CreateImplicit(SemaSYCLRef.getASTContext())); } return true; } @@ -1953,20 +1965,20 @@ class SyclKernelDecompMarker : public SyclKernelFieldHandler { // times. if (!FD->hasAttr()) FD->addAttr(SYCLRequiresDecompositionAttr::CreateImplicit( - SemaRef.getASTContext())); + SemaSYCLRef.getASTContext())); CollectionStack.back() = true; PointerStack.pop_back(); } else if (PointerStack.pop_back_val()) { if (!FD->hasAttr()) FD->addAttr( - SYCLGenerateNewTypeAttr::CreateImplicit(SemaRef.getASTContext())); + SYCLGenerateNewTypeAttr::CreateImplicit(SemaSYCLRef.getASTContext())); PointerStack.back() = true; } return true; } }; -static QualType ModifyAddressSpace(Sema &SemaRef, QualType Ty) { +static QualType ModifyAddressSpace(SemaSYCL &SemaSYCLRef, QualType Ty) { // USM allows to use raw pointers instead of buffers/accessors, but these // pointers point to the specially allocated memory. For pointer fields, // except for function pointer fields, we add a kernel argument with the @@ -1981,9 +1993,9 @@ static QualType ModifyAddressSpace(Sema &SemaRef, QualType Ty) { if (!PointeeTy->isFunctionType() && AS != LangAS::sycl_global_device && AS != LangAS::sycl_global_host) Quals.setAddressSpace(LangAS::sycl_global); - PointeeTy = SemaRef.getASTContext().getQualifiedType( + PointeeTy = SemaSYCLRef.getASTContext().getQualifiedType( PointeeTy.getUnqualifiedType(), Quals); - return SemaRef.getASTContext().getPointerType(PointeeTy); + return SemaSYCLRef.getASTContext().getPointerType(PointeeTy); } // This visitor is used to traverse a non-decomposed record/array to @@ -1996,7 +2008,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { IdentifierInfo *getModifiedName(IdentifierInfo *Id) { std::string Name = Id ? (Twine("__generated_") + Id->getName()).str() : "__generated_"; - return &SemaRef.getASTContext().Idents.get(Name); + return &SemaSYCLRef.getASTContext().Idents.get(Name); } // Create Decl for the new type we are generating. @@ -2004,7 +2016,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { // the visitor traverses kernel object record fields. void createNewType(const CXXRecordDecl *RD) { auto *ModifiedRD = CXXRecordDecl::Create( - SemaRef.getASTContext(), RD->getTagKind(), + SemaSYCLRef.getASTContext(), RD->getTagKind(), const_cast(RD->getDeclContext()), SourceLocation(), SourceLocation(), getModifiedName(RD->getIdentifier())); ModifiedRD->startDefinition(); @@ -2017,7 +2029,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { void addField(const FieldDecl *FD, QualType FieldTy) { assert(!ModifiedRecords.empty() && "ModifiedRecords should have at least 1 record"); - ASTContext &Ctx = SemaRef.getASTContext(); + ASTContext &Ctx = SemaSYCLRef.getASTContext(); auto *Field = FieldDecl::Create( Ctx, ModifiedRecords.back(), SourceLocation(), SourceLocation(), getModifiedName(FD->getIdentifier()), FieldTy, @@ -2032,9 +2044,9 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { void createBaseSpecifier(const CXXRecordDecl *Parent, const CXXRecordDecl *RD, const CXXBaseSpecifier &BS) { - TypeSourceInfo *TInfo = SemaRef.getASTContext().getTrivialTypeSourceInfo( + TypeSourceInfo *TInfo = SemaSYCLRef.getASTContext().getTrivialTypeSourceInfo( QualType(RD->getTypeForDecl(), 0), SourceLocation()); - CXXBaseSpecifier *ModifiedBase = SemaRef.CheckBaseSpecifier( + CXXBaseSpecifier *ModifiedBase = SemaSYCLRef.SemaRef.CheckBaseSpecifier( const_cast(Parent), SourceRange(), BS.isVirtual(), BS.getAccessSpecifier(), TInfo, SourceLocation()); ModifiedBases.push_back(ModifiedBase); @@ -2062,12 +2074,12 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { public: static constexpr const bool VisitInsideSimpleContainersWithPointer = true; static constexpr const bool VisitNthArrayElement = false; - SyclKernelPointerHandler(Sema &S, const CXXRecordDecl *RD) + SyclKernelPointerHandler(SemaSYCL &S, const CXXRecordDecl *RD) : SyclKernelFieldHandler(S) { createNewType(RD); } - SyclKernelPointerHandler(Sema &S) : SyclKernelFieldHandler(S) {} + SyclKernelPointerHandler(SemaSYCL &S) : SyclKernelFieldHandler(S) {} bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType Ty) final { createNewType(Ty->getAsCXXRecordDecl()); @@ -2107,9 +2119,9 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { QualType ModifiedArrayElement = ModifiedArrayElementsOrArray.pop_back_val(); const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(ArrayTy); + SemaSYCLRef.getASTContext().getAsConstantArrayType(ArrayTy); assert(CAT && "Should only be called on constant-size array."); - QualType ModifiedArray = SemaRef.getASTContext().getConstantArrayType( + QualType ModifiedArray = SemaSYCLRef.getASTContext().getConstantArrayType( ModifiedArrayElement, CAT->getSize(), const_cast(CAT->getSizeExpr()), CAT->getSizeModifier(), CAT->getIndexTypeCVRQualifiers()); @@ -2129,7 +2141,7 @@ class SyclKernelPointerHandler : public SyclKernelFieldHandler { } bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - QualType ModifiedPointerType = ModifyAddressSpace(SemaRef, FieldTy); + QualType ModifiedPointerType = ModifyAddressSpace(SemaSYCLRef, FieldTy); if (!isArrayElement(FD, FieldTy)) addField(FD, ModifiedPointerType); else @@ -2202,19 +2214,19 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // This only happens with the accessor types. StringRef Name = "_arg__base"; ParamDesc newParamDesc = - makeParamDesc(SemaRef.getASTContext(), Name, FieldTy); + makeParamDesc(SemaSYCLRef.getASTContext(), Name, FieldTy); addParam(newParamDesc, FieldTy); } // Add a parameter with specified name and type void addParam(StringRef Name, QualType ParamTy) { ParamDesc newParamDesc = - makeParamDesc(SemaRef.getASTContext(), Name, ParamTy); + makeParamDesc(SemaSYCLRef.getASTContext(), Name, ParamTy); addParam(newParamDesc, ParamTy); } void addParam(ParamDesc newParamDesc, QualType FieldTy) { // Create a new ParmVarDecl based on the new info. - ASTContext &Ctx = SemaRef.getASTContext(); + ASTContext &Ctx = SemaSYCLRef.getASTContext(); auto *NewParam = ParmVarDecl::Create( Ctx, KernelDecl, SourceLocation(), SourceLocation(), std::get<1>(newParamDesc), std::get<0>(newParamDesc), @@ -2253,7 +2265,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { void handleNoAliasProperty(ParmVarDecl *Param, QualType PropTy, SourceLocation Loc) { - ASTContext &Ctx = SemaRef.getASTContext(); + ASTContext &Ctx = SemaSYCLRef.getASTContext(); Param->addAttr(RestrictAttr::CreateImplicit(Ctx, Loc)); } @@ -2264,11 +2276,11 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // If we have more than 1 buffer_location properties on a single // accessor - emit an error if (Param->hasAttr()) { - SemaRef.Diag(Loc, diag::err_sycl_compiletime_property_duplication) + SemaSYCLRef.Diag(Loc, diag::err_sycl_compiletime_property_duplication) << "buffer_location"; return; } - ASTContext &Ctx = SemaRef.getASTContext(); + ASTContext &Ctx = SemaSYCLRef.getASTContext(); const auto *PropDecl = cast(PropTy->getAsRecordDecl()); const auto BufferLoc = PropDecl->getTemplateArgs()[0]; @@ -2283,7 +2295,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { handleAccessorPropertyList(Params.back(), RecordDecl, Loc); // If "accessor" type check if read only - if (Sema::isSyclType(FieldTy, SYCLTypeAttr::accessor)) { + if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::accessor)) { // Get access mode of accessor. const auto *AccessorSpecializationDecl = cast(RecordDecl); @@ -2291,13 +2303,13 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { AccessorSpecializationDecl->getTemplateArgs().get(2); if (isReadOnlyAccessor(AccessModeArg)) Params.back()->addAttr( - SYCLAccessorReadonlyAttr::CreateImplicit(SemaRef.getASTContext())); + SYCLAccessorReadonlyAttr::CreateImplicit(SemaSYCLRef.getASTContext())); } // Add implicit attribute to parameter decl when it is a read only // SYCL accessor. Params.back()->addAttr( - SYCLAccessorPtrAttr::CreateImplicit(SemaRef.getASTContext())); + SYCLAccessorPtrAttr::CreateImplicit(SemaSYCLRef.getASTContext())); } // All special SYCL objects must have __init method. We extract types for @@ -2324,7 +2336,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Propagate add_ir_attributes_kernel_parameter attribute. if (const auto *AddIRAttr = Param->getAttr()) - Params.back()->addAttr(AddIRAttr->clone(SemaRef.getASTContext())); + Params.back()->addAttr(AddIRAttr->clone(SemaSYCLRef.getASTContext())); // FIXME: This code is temporary, and will be removed once __init_esimd // is removed and property list refactored. @@ -2374,8 +2386,8 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // pointers in 'right' address space. PointerHandler.getNewRecordType() // returns this generated type. QualType GenerateNewRecordType(const CXXRecordDecl *RD) { - SyclKernelPointerHandler PointerHandler(SemaRef, RD); - KernelObjVisitor Visitor{SemaRef}; + SyclKernelPointerHandler PointerHandler(SemaSYCLRef, RD); + KernelObjVisitor Visitor{SemaSYCLRef}; Visitor.VisitRecordBases(RD, PointerHandler); Visitor.VisitRecordFields(RD, PointerHandler); return PointerHandler.getNewRecordType(); @@ -2387,29 +2399,29 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // a new array with all pointers in the required address space. QualType GenerateNewArrayType(FieldDecl *FD, QualType FieldTy) { const auto *Owner = dyn_cast(FD->getParent()); - SyclKernelPointerHandler PointerHandler(SemaRef); - KernelObjVisitor Visitor{SemaRef}; + SyclKernelPointerHandler PointerHandler(SemaSYCLRef); + KernelObjVisitor Visitor{SemaSYCLRef}; Visitor.visitArray(Owner, FD, FieldTy, PointerHandler); return PointerHandler.getNewArrayType(); } public: static constexpr const bool VisitInsideSimpleContainers = false; - SyclKernelDeclCreator(Sema &S, SourceLocation Loc, bool IsInline, + SyclKernelDeclCreator(SemaSYCL &S, SourceLocation Loc, bool IsInline, bool IsSIMDKernel, FunctionDecl *SYCLKernel) : SyclKernelFieldHandler(S), KernelDecl( createKernelDecl(S.getASTContext(), Loc, IsInline, IsSIMDKernel)), - FuncContext(SemaRef, KernelDecl) { + FuncContext(SemaSYCLRef.SemaRef, KernelDecl) { S.addSyclOpenCLKernel(SYCLKernel, KernelDecl); if (const auto *AddIRAttrFunc = SYCLKernel->getAttr()) - KernelDecl->addAttr(AddIRAttrFunc->clone(SemaRef.getASTContext())); + KernelDecl->addAttr(AddIRAttrFunc->clone(SemaSYCLRef.getASTContext())); } ~SyclKernelDeclCreator() { - ASTContext &Ctx = SemaRef.getASTContext(); + ASTContext &Ctx = SemaSYCLRef.getASTContext(); FunctionProtoType::ExtProtoInfo Info(CC_OpenCLKernel); SmallVector ArgTys; @@ -2426,9 +2438,9 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // to TransformStmt in replaceWithLocalClone can diagnose something that got // diagnosed on the actual kernel. KernelDecl->addAttr( - SYCLKernelAttr::CreateImplicit(SemaRef.getASTContext())); + SYCLKernelAttr::CreateImplicit(SemaSYCLRef.getASTContext())); - SemaRef.addSyclDeviceDecl(KernelDecl); + SemaSYCLRef.addSyclDeviceDecl(KernelDecl); } bool enterStruct(const CXXRecordDecl *, FieldDecl *, QualType) final { @@ -2489,12 +2501,12 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { RecordDecl *wrapField(FieldDecl *Field, QualType FieldTy) { RecordDecl *WrapperClass = - SemaRef.getASTContext().buildImplicitRecord("__wrapper_class"); + SemaSYCLRef.getASTContext().buildImplicitRecord("__wrapper_class"); WrapperClass->startDefinition(); Field = FieldDecl::Create( - SemaRef.getASTContext(), WrapperClass, SourceLocation(), + SemaSYCLRef.getASTContext(), WrapperClass, SourceLocation(), SourceLocation(), /*Id=*/nullptr, FieldTy, - SemaRef.getASTContext().getTrivialTypeSourceInfo(FieldTy, + SemaSYCLRef.getASTContext().getTrivialTypeSourceInfo(FieldTy, SourceLocation()), /*BW=*/nullptr, /*Mutable=*/false, /*InitStyle=*/ICIS_NoInit); Field->setAccess(AS_public); @@ -2504,7 +2516,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { }; bool handlePointerType(FieldDecl *FD, QualType FieldTy) final { - QualType ModTy = ModifyAddressSpace(SemaRef, FieldTy); + QualType ModTy = ModifyAddressSpace(SemaSYCLRef, FieldTy); // When the kernel is generated, struct type kernel arguments are // decomposed; i.e. the parameters of the kernel are the fields of the // struct, and not the struct itself. This causes an error in the backend @@ -2513,7 +2525,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // struct are wrapped in a generated '__wrapper_class'. if (StructDepth) { RecordDecl *WrappedPointer = wrapField(FD, ModTy); - ModTy = SemaRef.getASTContext().getRecordType(WrappedPointer); + ModTy = SemaSYCLRef.getASTContext().getRecordType(WrappedPointer); } addParam(FD, ModTy); @@ -2529,7 +2541,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Arrays are wrapped in a struct since they cannot be passed directly. RecordDecl *WrappedArray = wrapField(FD, ArrayTy); - addParam(FD, SemaRef.getASTContext().getRecordType(WrappedArray)); + addParam(FD, SemaSYCLRef.getASTContext().getRecordType(WrappedArray)); return true; } @@ -2572,7 +2584,7 @@ class SyclKernelDeclCreator : public SyclKernelFieldHandler { // Generate kernel argument to initialize specialization constants. void handleSyclKernelHandlerType() { - ASTContext &Context = SemaRef.getASTContext(); + ASTContext &Context = SemaSYCLRef.getASTContext(); StringRef Name = "_arg__specialization_constants_buffer"; addParam(Name, Context.getPointerType(Context.getAddrSpaceQualType( Context.CharTy, LangAS::sycl_global))); @@ -2661,14 +2673,14 @@ class ESIMDKernelDiagnostics : public SyclKernelFieldHandler { const CXXRecordDecl *RecordDecl = FieldTy->getAsCXXRecordDecl(); if (IsESIMD && !isSyclAccessorType(FieldTy)) - return SemaRef.Diag(KernelLoc, + return SemaSYCLRef.Diag(KernelLoc, diag::err_sycl_esimd_not_supported_for_type) << RecordDecl; return true; } public: - ESIMDKernelDiagnostics(Sema &S, SourceLocation Loc, bool IsESIMD) + ESIMDKernelDiagnostics(SemaSYCL &S, SourceLocation Loc, bool IsESIMD) : SyclKernelFieldHandler(S), KernelLoc(Loc), IsESIMD(IsESIMD) {} bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final { @@ -2688,7 +2700,7 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { void addParam(QualType ArgTy) { SizeOfParams += - SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); + SemaSYCLRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); } bool handleSpecialType(QualType FieldTy) { @@ -2706,12 +2718,12 @@ class SyclKernelArgsSizeChecker : public SyclKernelFieldHandler { public: static constexpr const bool VisitInsideSimpleContainers = false; - SyclKernelArgsSizeChecker(Sema &S, SourceLocation Loc, bool IsESIMD) + SyclKernelArgsSizeChecker(SemaSYCL &S, SourceLocation Loc, bool IsESIMD) : SyclKernelFieldHandler(S), KernelLoc(Loc), IsESIMD(IsESIMD) {} ~SyclKernelArgsSizeChecker() { if (SizeOfParams > MaxKernelArgsSize) - SemaRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_big_args) + SemaSYCLRef.Diag(KernelLoc, diag::warn_sycl_kernel_too_big_args) << SizeOfParams << MaxKernelArgsSize; } @@ -2776,9 +2788,9 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { NameToEmitInDescription = KernelArgParent->getName(); unsigned KernelArgSize = - SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + SemaSYCLRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); - SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( + SemaSYCLRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( DC.getKernelDecl(), NameToEmitInDescription, IsCompilerGeneratedType ? "Compiler generated" : KernelArgType.getAsString(), @@ -2803,8 +2815,8 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { StringRef KernelArgDescription, bool IsCompilerGeneratedType = false) { unsigned KernelArgSize = - SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); - SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( + SemaSYCLRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + SemaSYCLRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( DC.getKernelDecl(), KernelArgType.getAsString(), IsCompilerGeneratedType ? "Compiler generated" : KernelArgType.getAsString(), @@ -2815,8 +2827,8 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { // Handles specialization constants. void addParam(QualType KernelArgType, std::string KernelArgDescription) { unsigned KernelArgSize = - SemaRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); - SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( + SemaSYCLRef.getASTContext().getTypeSizeInChars(KernelArgType).getQuantity(); + SemaSYCLRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( DC.getKernelDecl(), "", KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDescription), ""); @@ -2824,7 +2836,7 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { public: static constexpr const bool VisitInsideSimpleContainers = false; - SyclOptReportCreator(Sema &S, SyclKernelDeclCreator &DC, SourceLocation Loc) + SyclOptReportCreator(SemaSYCL &S, SyclKernelDeclCreator &DC, SourceLocation Loc) : SyclKernelFieldHandler(S), DC(DC), KernelInvocationLoc(Loc) {} bool handleSyclSpecialType(FieldDecl *FD, QualType FieldTy) final { @@ -2838,10 +2850,10 @@ class SyclOptReportCreator : public SyclKernelFieldHandler { std::string KernelArgDescription = "base class " + FieldTy.getAsString(); for (const auto *Param : DC.getParamVarDeclsForCurrentField()) { QualType KernelArgType = Param->getType(); - unsigned KernelArgSize = SemaRef.getASTContext() + unsigned KernelArgSize = SemaSYCLRef.getASTContext() .getTypeSizeInChars(KernelArgType) .getQuantity(); - SemaRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( + SemaSYCLRef.getDiagnostics().getSYCLOptReport().AddKernelArgs( DC.getKernelDecl(), FieldTy.getAsString(), KernelArgType.getAsString(), KernelInvocationLoc, KernelArgSize, getKernelArgDesc(KernelArgDescription), ""); @@ -2948,7 +2960,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { LocalClone->setIsUsed(); std::pair MappingPair = std::make_pair(OriginalParam, LocalClone); - KernelBodyTransform KBT(MappingPair, SemaRef); + KernelBodyTransform KBT(MappingPair, SemaSYCLRef.SemaRef); return KBT.TransformStmt(FunctionBody).get(); } @@ -2957,7 +2969,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // statements in advance to allocate it, so we cannot do this as we go along. CompoundStmt *createKernelBody() { // Push the Kernel function scope to ensure the scope isn't empty - SemaRef.PushFunctionScope(); + SemaSYCLRef.SemaRef.PushFunctionScope(); // Initialize kernel object local clone assert(CollectionInitExprs.size() == 1 && @@ -2984,7 +2996,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { BodyStmts.insert(BodyStmts.end(), FinalizeStmts.begin(), FinalizeStmts.end()); - return CompoundStmt::Create(SemaRef.getASTContext(), BodyStmts, + return CompoundStmt::Create(SemaSYCLRef.getASTContext(), BodyStmts, FPOptionsOverride(), {}, {}); } @@ -2996,7 +3008,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Mark kernel object with work-group scope attribute to avoid work-item // scope memory allocation. KernelObjClone->addAttr(SYCLScopeAttr::CreateImplicit( - SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); + SemaSYCLRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); assert(CallOperator && "non callable object is passed as kernel obj"); // Mark the function that it "works" in a work group scope: @@ -3010,9 +3022,9 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // SYCL headers. if (!CallOperator->hasAttr()) { CallOperator->addAttr(SYCLScopeAttr::CreateImplicit( - SemaRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); + SemaSYCLRef.getASTContext(), SYCLScopeAttr::Level::WorkGroup)); // Search and mark wait_for calls: - MarkWIScopeFnVisitor MarkWIScope(SemaRef.getASTContext()); + MarkWIScopeFnVisitor MarkWIScope(SemaSYCLRef.getASTContext()); MarkWIScope.TraverseDecl(CallOperator); // Now mark local variables declared in the PFWG lambda with work group // scope attribute @@ -3026,8 +3038,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DeclCreator.getParamVarDeclsForCurrentField()[0]; QualType ParamType = KernelParameter->getOriginalType(); - Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, - KernelCallerSrcLoc); + Expr *DRE = SemaSYCLRef.SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, + VK_LValue, KernelCallerSrcLoc); return DRE; } @@ -3038,8 +3050,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DeclCreator.getParamVarDeclsForCurrentField()[0]; QualType ParamType = KernelParameter->getOriginalType(); - Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, - KernelCallerSrcLoc); + Expr *DRE = SemaSYCLRef.SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, + VK_LValue, KernelCallerSrcLoc); // Struct Type kernel arguments are decomposed. The pointer fields are // then wrapped inside a compiler generated struct. Therefore when @@ -3052,13 +3064,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ParamType = Pointer->getType(); } - DRE = ImplicitCastExpr::Create(SemaRef.Context, ParamType, + DRE = ImplicitCastExpr::Create(SemaSYCLRef.getASTContext(), ParamType, CK_LValueToRValue, DRE, /*BasePath=*/nullptr, VK_PRValue, FPOptionsOverride()); if (PointerTy->getPointeeType().getAddressSpace() != ParamType->getPointeeType().getAddressSpace()) - DRE = ImplicitCastExpr::Create(SemaRef.Context, PointerTy, + DRE = ImplicitCastExpr::Create(SemaSYCLRef.getASTContext(), PointerTy, CK_AddressSpaceConversion, DRE, nullptr, VK_PRValue, FPOptionsOverride()); @@ -3069,8 +3081,8 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { ParmVarDecl *KernelParameter = DeclCreator.getParamVarDeclsForCurrentField()[0]; QualType ParamType = KernelParameter->getOriginalType(); - Expr *DRE = SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, VK_LValue, - KernelCallerSrcLoc); + Expr *DRE = SemaSYCLRef.SemaRef.BuildDeclRefExpr(KernelParameter, ParamType, + VK_LValue, KernelCallerSrcLoc); // Unwrap the array. CXXRecordDecl *WrapperStruct = ParamType->getAsCXXRecordDecl(); @@ -3083,7 +3095,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // array, returns an element initializer. InitializedEntity getFieldEntity(FieldDecl *FD, QualType Ty) { if (isArrayElement(FD, Ty)) - return InitializedEntity::InitializeElement(SemaRef.getASTContext(), + return InitializedEntity::InitializeElement(SemaSYCLRef.getASTContext(), ArrayInfos.back().second, ArrayInfos.back().first); return InitializedEntity::InitializeMember(FD, &VarEntity); @@ -3102,35 +3114,37 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { void addFieldInit(FieldDecl *FD, QualType Ty, MultiExprArg ParamRef, InitializationKind InitKind, InitializedEntity Entity) { - InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef); - ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef); + InitializationSequence InitSeq(SemaSYCLRef.SemaRef, Entity, InitKind, ParamRef); + ExprResult Init = InitSeq.Perform(SemaSYCLRef.SemaRef, Entity, InitKind, ParamRef); InitListExpr *ParentILE = CollectionInitExprs.back(); - ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + ParentILE->updateInit(SemaSYCLRef.getASTContext(), ParentILE->getNumInits(), Init.get()); } void addBaseInit(const CXXBaseSpecifier &BS, QualType Ty, InitializationKind InitKind) { InitializedEntity Entity = InitializedEntity::InitializeBase( - SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, std::nullopt); - ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, std::nullopt); + SemaSYCLRef.getASTContext(), &BS, /*IsInheritedVirtualBase*/ false, + &VarEntity); + InitializationSequence InitSeq(SemaSYCLRef.SemaRef, Entity, InitKind, std::nullopt); + ExprResult Init = InitSeq.Perform(SemaSYCLRef.SemaRef, Entity, InitKind, std::nullopt); InitListExpr *ParentILE = CollectionInitExprs.back(); - ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + ParentILE->updateInit(SemaSYCLRef.getASTContext(), ParentILE->getNumInits(), Init.get()); } void addBaseInit(const CXXBaseSpecifier &BS, QualType Ty, InitializationKind InitKind, MultiExprArg Args) { InitializedEntity Entity = InitializedEntity::InitializeBase( - SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, Args); - ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, Args); + SemaSYCLRef.getASTContext(), &BS, /*IsInheritedVirtualBase*/ false, + &VarEntity); + InitializationSequence InitSeq(SemaSYCLRef.SemaRef, Entity, InitKind, Args); + ExprResult Init = InitSeq.Perform(SemaSYCLRef.SemaRef, Entity, InitKind, Args); InitListExpr *ParentILE = CollectionInitExprs.back(); - ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + ParentILE->updateInit(SemaSYCLRef.getASTContext(), ParentILE->getNumInits(), Init.get()); } @@ -3139,14 +3153,15 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc); InitializedEntity Entity = InitializedEntity::InitializeBase( - SemaRef.Context, &BS, /*IsInheritedVirtualBase*/ false, &VarEntity); + SemaSYCLRef.getASTContext(), &BS, /*IsInheritedVirtualBase*/ false, + &VarEntity); Expr *ParamRef = createParamReferenceExpr(); - InitializationSequence InitSeq(SemaRef, Entity, InitKind, ParamRef); - ExprResult Init = InitSeq.Perform(SemaRef, Entity, InitKind, ParamRef); + InitializationSequence InitSeq(SemaSYCLRef.SemaRef, Entity, InitKind, ParamRef); + ExprResult Init = InitSeq.Perform(SemaSYCLRef.SemaRef, Entity, InitKind, ParamRef); InitListExpr *ParentILE = CollectionInitExprs.back(); - ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + ParentILE->updateInit(SemaSYCLRef.getASTContext(), ParentILE->getNumInits(), Init.get()); } @@ -3157,23 +3172,24 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { } Expr *createGetAddressOf(Expr *E) { - return UnaryOperator::Create(SemaRef.Context, E, UO_AddrOf, - SemaRef.Context.getPointerType(E->getType()), - VK_PRValue, OK_Ordinary, KernelCallerSrcLoc, - false, SemaRef.CurFPFeatureOverrides()); + return UnaryOperator::Create( + SemaSYCLRef.getASTContext(), E, UO_AddrOf, + SemaSYCLRef.getASTContext().getPointerType(E->getType()), VK_PRValue, + OK_Ordinary, KernelCallerSrcLoc, false, + SemaSYCLRef.SemaRef.CurFPFeatureOverrides()); } Expr *createDerefOp(Expr *E) { - return UnaryOperator::Create(SemaRef.Context, E, UO_Deref, + return UnaryOperator::Create(SemaSYCLRef.getASTContext(), E, UO_Deref, E->getType()->getPointeeType(), VK_LValue, OK_Ordinary, KernelCallerSrcLoc, false, - SemaRef.CurFPFeatureOverrides()); + SemaSYCLRef.SemaRef.CurFPFeatureOverrides()); } Expr *createReinterpretCastExpr(Expr *E, QualType To) { return CXXReinterpretCastExpr::Create( - SemaRef.Context, To, VK_PRValue, CK_BitCast, E, - /*Path=*/nullptr, SemaRef.Context.getTrivialTypeSourceInfo(To), + SemaSYCLRef.getASTContext(), To, VK_PRValue, CK_BitCast, E, + /*Path=*/nullptr, SemaSYCLRef.getASTContext().getTrivialTypeSourceInfo(To), SourceLocation(), SourceLocation(), SourceRange()); } @@ -3185,7 +3201,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { Expr *RCE = createReinterpretCastExpr( createGetAddressOf(createParamReferenceExpr()), - SemaRef.Context.getPointerType(Ty)); + SemaSYCLRef.getASTContext().getPointerType(Ty)); Expr *Initializer = createDerefOp(RCE); addFieldInit(FD, Ty, Initializer); } @@ -3198,7 +3214,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // } Expr *RCE = createReinterpretCastExpr( createGetAddressOf(createParamReferenceExpr()), - SemaRef.Context.getPointerType(Ty)); + SemaSYCLRef.getASTContext().getPointerType(Ty)); Expr *Initializer = createDerefOp(RCE); InitializationKind InitKind = InitializationKind::CreateCopy(KernelCallerSrcLoc, KernelCallerSrcLoc); @@ -3207,7 +3223,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { MemberExpr *buildMemberExpr(Expr *Base, ValueDecl *Member) { DeclAccessPair MemberDAP = DeclAccessPair::make(Member, AS_none); - MemberExpr *Result = SemaRef.BuildMemberExpr( + MemberExpr *Result = SemaSYCLRef.SemaRef.BuildMemberExpr( Base, /*IsArrow */ false, KernelCallerSrcLoc, NestedNameSpecifierLoc(), KernelCallerSrcLoc, Member, MemberDAP, /*HadMultipleCandidates*/ false, @@ -3238,23 +3254,23 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { DeclCreator.getParamVarDeclsForCurrentField(); for (size_t I = 0; I < NumParams; ++I) { QualType ParamType = KernelParameters[I]->getOriginalType(); - ParamDREs[I] = SemaRef.BuildDeclRefExpr(KernelParameters[I], ParamType, - VK_LValue, KernelCallerSrcLoc); + ParamDREs[I] = SemaSYCLRef.SemaRef.BuildDeclRefExpr( + KernelParameters[I], ParamType, VK_LValue, KernelCallerSrcLoc); } MemberExpr *MethodME = buildMemberExpr(MemberExprBases.back(), Method); QualType ResultTy = Method->getReturnType(); ExprValueKind VK = Expr::getValueKindForType(ResultTy); - ResultTy = ResultTy.getNonLValueExprType(SemaRef.Context); + ResultTy = ResultTy.getNonLValueExprType(SemaSYCLRef.getASTContext()); llvm::SmallVector ParamStmts; const auto *Proto = cast(Method->getType()); - SemaRef.GatherArgumentsForCall(KernelCallerSrcLoc, Method, Proto, 0, - ParamDREs, ParamStmts); + SemaSYCLRef.SemaRef.GatherArgumentsForCall(KernelCallerSrcLoc, Method, Proto, 0, + ParamDREs, ParamStmts); // [kernel_obj or wrapper object].accessor.__init(_ValueType*, // range, range, id) AddTo.push_back(CXXMemberCallExpr::Create( - SemaRef.Context, MethodME, ParamStmts, ResultTy, VK, KernelCallerSrcLoc, + SemaSYCLRef.getASTContext(), MethodME, ParamStmts, ResultTy, VK, KernelCallerSrcLoc, FPOptionsOverride())); } @@ -3262,22 +3278,22 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // of this to append into. void addCollectionInitListExpr(const CXXRecordDecl *RD) { const ASTRecordLayout &Info = - SemaRef.getASTContext().getASTRecordLayout(RD); + SemaSYCLRef.getASTContext().getASTRecordLayout(RD); uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases(); addCollectionInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); } InitListExpr *createInitListExpr(const CXXRecordDecl *RD) { const ASTRecordLayout &Info = - SemaRef.getASTContext().getASTRecordLayout(RD); + SemaSYCLRef.getASTContext().getASTRecordLayout(RD); uint64_t NumInitExprs = Info.getFieldCount() + RD->getNumBases(); return createInitListExpr(QualType(RD->getTypeForDecl(), 0), NumInitExprs); } InitListExpr *createInitListExpr(QualType InitTy, uint64_t NumChildInits) { - InitListExpr *ILE = new (SemaRef.getASTContext()) InitListExpr( - SemaRef.getASTContext(), KernelCallerSrcLoc, {}, KernelCallerSrcLoc); - ILE->reserveInits(SemaRef.getASTContext(), NumChildInits); + InitListExpr *ILE = new (SemaSYCLRef.getASTContext()) InitListExpr( + SemaSYCLRef.getASTContext(), KernelCallerSrcLoc, {}, KernelCallerSrcLoc); + ILE->reserveInits(SemaSYCLRef.getASTContext(), NumChildInits); ILE->setType(InitTy); return ILE; @@ -3289,7 +3305,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitListExpr *ILE = createInitListExpr(InitTy, NumChildInits); InitListExpr *ParentILE = CollectionInitExprs.back(); - ParentILE->updateInit(SemaRef.getASTContext(), ParentILE->getNumInits(), + ParentILE->updateInit(SemaSYCLRef.getASTContext(), ParentILE->getNumInits(), ILE); CollectionInitExprs.push_back(ILE); @@ -3343,7 +3359,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Generate __init call for kernel handler argument void handleSpecialType(QualType KernelHandlerTy) { DeclRefExpr *KernelHandlerCloneRef = - DeclRefExpr::Create(SemaRef.Context, NestedNameSpecifierLoc(), + DeclRefExpr::Create(SemaSYCLRef.getASTContext(), NestedNameSpecifierLoc(), KernelCallerSrcLoc, KernelHandlerClone, false, DeclarationNameInfo(), KernelHandlerTy, VK_LValue); const auto *RecordDecl = @@ -3366,21 +3382,23 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { InitializedEntity::InitializeVariable(KernelHandlerClone); InitializationKind InitKind = InitializationKind::CreateDefault(KernelCallerSrcLoc); - InitializationSequence InitSeq(SemaRef, VarEntity, InitKind, std::nullopt); - ExprResult Init = InitSeq.Perform(SemaRef, VarEntity, InitKind, std::nullopt); + InitializationSequence InitSeq(SemaSYCLRef.SemaRef, VarEntity, InitKind, + std::nullopt); + ExprResult Init = + InitSeq.Perform(SemaSYCLRef.SemaRef, VarEntity, InitKind, std::nullopt); KernelHandlerClone->setInit( - SemaRef.MaybeCreateExprWithCleanups(Init.get())); + SemaSYCLRef.SemaRef.MaybeCreateExprWithCleanups(Init.get())); KernelHandlerClone->setInitStyle(VarDecl::CallInit); } Expr *createArraySubscriptExpr(uint64_t Index, Expr *ArrayRef) { - QualType SizeT = SemaRef.getASTContext().getSizeType(); + QualType SizeT = SemaSYCLRef.getASTContext().getSizeType(); llvm::APInt IndexVal{ - static_cast(SemaRef.getASTContext().getTypeSize(SizeT)), + static_cast(SemaSYCLRef.getASTContext().getTypeSize(SizeT)), Index, SizeT->isSignedIntegerType()}; auto IndexLiteral = IntegerLiteral::Create( - SemaRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc); - ExprResult IndexExpr = SemaRef.CreateBuiltinArraySubscriptExpr( + SemaSYCLRef.getASTContext(), IndexVal, SizeT, KernelCallerSrcLoc); + ExprResult IndexExpr = SemaSYCLRef.SemaRef.CreateBuiltinArraySubscriptExpr( ArrayRef, KernelCallerSrcLoc, IndexLiteral, KernelCallerSrcLoc); assert(!IndexExpr.isInvalid()); return IndexExpr.get(); @@ -3399,7 +3417,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { void addArrayElementInit(FieldDecl *FD, QualType T) { Expr *RCE = createReinterpretCastExpr( createGetAddressOf(ArrayParamBases.pop_back_val()), - SemaRef.Context.getPointerType(T)); + SemaSYCLRef.getASTContext().getPointerType(T)); Expr *Initializer = createDerefOp(RCE); addFieldInit(FD, T, Initializer); } @@ -3411,7 +3429,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // generate required array subscript expressions. void createArrayInit(FieldDecl *FD, QualType T) { const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(T); + SemaSYCLRef.getASTContext().getAsConstantArrayType(T); if (!CAT) { addArrayElementInit(FD, T); @@ -3452,7 +3470,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { public: static constexpr const bool VisitInsideSimpleContainers = false; - SyclKernelBodyCreator(Sema &S, SyclKernelDeclCreator &DC, + SyclKernelBodyCreator(SemaSYCL &S, SyclKernelDeclCreator &DC, const CXXRecordDecl *KernelObj, FunctionDecl *KernelCallerFunc, bool IsSIMDKernel, CXXMethodDecl *CallOperator) @@ -3466,13 +3484,13 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CollectionInitExprs.push_back(createInitListExpr(KernelObj)); annotateHierarchicalParallelismAPICalls(); - Stmt *DS = new (S.Context) DeclStmt(DeclGroupRef(KernelObjClone), - KernelCallerSrcLoc, KernelCallerSrcLoc); + Stmt *DS = new (S.getASTContext()) DeclStmt( + DeclGroupRef(KernelObjClone), KernelCallerSrcLoc, KernelCallerSrcLoc); BodyStmts.push_back(DS); DeclRefExpr *KernelObjCloneRef = DeclRefExpr::Create( - S.Context, NestedNameSpecifierLoc(), KernelCallerSrcLoc, KernelObjClone, - false, DeclarationNameInfo(), QualType(KernelObj->getTypeForDecl(), 0), - VK_LValue); + S.getASTContext(), NestedNameSpecifierLoc(), KernelCallerSrcLoc, + KernelObjClone, false, DeclarationNameInfo(), + QualType(KernelObj->getTypeForDecl(), 0), VK_LValue); MemberExprBases.push_back(KernelObjCloneRef); } @@ -3541,12 +3559,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { void handleSyclKernelHandlerType(ParmVarDecl *KernelHandlerArg) { // Create and default initialize local clone of kernel handler - createKernelHandlerClone(SemaRef.getASTContext(), + createKernelHandlerClone(SemaSYCLRef.getASTContext(), DeclCreator.getKernelDecl(), KernelHandlerArg); // Add declaration statement to openCL kernel body Stmt *DS = - new (SemaRef.Context) DeclStmt(DeclGroupRef(KernelHandlerClone), + new (SemaSYCLRef.getASTContext()) DeclStmt(DeclGroupRef(KernelHandlerClone), KernelCallerSrcLoc, KernelCallerSrcLoc); BodyStmts.push_back(DS); @@ -3555,7 +3573,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // call if target does not have native support for specialization constants. // Here, specialization_constants_buffer is the compiler generated kernel // argument of type char*. - if (!isDefaultSPIRArch(SemaRef.Context)) + if (!isDefaultSPIRArch(SemaSYCLRef.getASTContext())) handleSpecialType(KernelHandlerArg->getType()); } @@ -3582,11 +3600,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { CXXCastPath BasePath; QualType DerivedTy(RD->getTypeForDecl(), 0); QualType BaseTy = BS.getType(); - SemaRef.CheckDerivedToBaseConversion(DerivedTy, BaseTy, KernelCallerSrcLoc, - SourceRange(), &BasePath, - /*IgnoreBaseAccess*/ true); + SemaSYCLRef.SemaRef.CheckDerivedToBaseConversion( + DerivedTy, BaseTy, KernelCallerSrcLoc, SourceRange(), &BasePath, + /*IgnoreBaseAccess*/ true); auto Cast = ImplicitCastExpr::Create( - SemaRef.Context, BaseTy, CK_DerivedToBase, MemberExprBases.back(), + SemaSYCLRef.getASTContext(), BaseTy, CK_DerivedToBase, + MemberExprBases.back(), /* CXXCastPath=*/&BasePath, VK_LValue, FPOptionsOverride()); MemberExprBases.push_back(Cast); addCollectionInitListExpr(BaseTy->getAsCXXRecordDecl()); @@ -3604,7 +3623,7 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { bool enterArray(FieldDecl *FD, QualType ArrayType, QualType ElementType) final { const ConstantArrayType *CAT = - SemaRef.getASTContext().getAsConstantArrayType(ArrayType); + SemaSYCLRef.getASTContext().getAsConstantArrayType(ArrayType); assert(CAT && "Should only be called on constant-size array."); uint64_t ArraySize = CAT->getSize().getZExtValue(); addCollectionInitListExpr(ArrayType, ArraySize); @@ -3648,12 +3667,12 @@ class SyclKernelBodyCreator : public SyclKernelFieldHandler { // Kernels are only the unnamed-lambda feature if the feature is enabled, AND // the first template argument has been corrected by the library to match the // functor type. -static bool IsSYCLUnnamedKernel(Sema &SemaRef, const FunctionDecl *FD) { - if (!SemaRef.getLangOpts().SYCLUnnamedLambda) +static bool IsSYCLUnnamedKernel(SemaSYCL &SemaSYCLRef, const FunctionDecl *FD) { + if (!SemaSYCLRef.getLangOpts().SYCLUnnamedLambda) return false; QualType FunctorTy = GetSYCLKernelObjectType(FD); - QualType TmplArgTy = calculateKernelNameType(SemaRef.Context, FD); - return SemaRef.Context.hasSameType(FunctorTy, TmplArgTy); + QualType TmplArgTy = calculateKernelNameType(SemaSYCLRef.getASTContext(), FD); + return SemaSYCLRef.getASTContext().hasSameType(FunctorTy, TmplArgTy); } class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { @@ -3666,12 +3685,12 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { int64_t offsetOf(const FieldDecl *FD, QualType ArgTy) const { return isArrayElement(FD, ArgTy) ? 0 - : SemaRef.getASTContext().getFieldOffset(FD) / 8; + : SemaSYCLRef.getASTContext().getFieldOffset(FD) / 8; } int64_t offsetOf(const CXXRecordDecl *RD, const CXXRecordDecl *Base) const { const ASTRecordLayout &Layout = - SemaRef.getASTContext().getASTRecordLayout(RD); + SemaSYCLRef.getASTContext().getASTRecordLayout(RD); return Layout.getBaseClassOffset(Base).getQuantity(); } @@ -3682,20 +3701,20 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { void addParam(QualType ArgTy, SYCLIntegrationHeader::kernel_param_kind_t Kind, uint64_t OffsetAdj) { uint64_t Size; - Size = SemaRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); + Size = SemaSYCLRef.getASTContext().getTypeSizeInChars(ArgTy).getQuantity(); Header.addParamDesc(Kind, static_cast(Size), static_cast(CurOffset + OffsetAdj)); } public: static constexpr const bool VisitInsideSimpleContainers = false; - SyclKernelIntHeaderCreator(bool IsESIMD, Sema &S, SYCLIntegrationHeader &H, + SyclKernelIntHeaderCreator(bool IsESIMD, SemaSYCL &S, SYCLIntegrationHeader &H, const CXXRecordDecl *KernelObj, QualType NameType, FunctionDecl *KernelFunc) : SyclKernelFieldHandler(S), Header(H) { // The header needs to access the kernel object size. - int64_t ObjSize = SemaRef.getASTContext() + int64_t ObjSize = SemaSYCLRef.getASTContext() .getTypeSizeInChars(KernelObj->getTypeForDecl()) .getQuantity(); Header.startKernel(KernelFunc, NameType, KernelObj->getLocation(), IsESIMD, @@ -3732,18 +3751,18 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { Header.addParamDesc(SYCLIntegrationHeader::kind_accessor, Info, CurOffset + offsetOf(FD, FieldTy)); - } else if (Sema::isSyclType(FieldTy, SYCLTypeAttr::stream)) { + } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::stream)) { addParam(FD, FieldTy, SYCLIntegrationHeader::kind_stream); - } else if (Sema::isSyclType(FieldTy, SYCLTypeAttr::sampler) || - Sema::isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) || - Sema::isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) { + } else if (SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::sampler) || + SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_ptr) || + SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::annotated_arg)) { CXXMethodDecl *InitMethod = getMethodByName(ClassTy, InitMethodName); assert(InitMethod && "type must have __init method"); const ParmVarDecl *InitArg = InitMethod->getParamDecl(0); assert(InitArg && "Init method must have arguments"); QualType T = InitArg->getType(); SYCLIntegrationHeader::kernel_param_kind_t ParamKind = - Sema::isSyclType(FieldTy, SYCLTypeAttr::sampler) + SemaSYCL::isSyclType(FieldTy, SYCLTypeAttr::sampler) ? SYCLIntegrationHeader::kind_sampler : (T->isPointerType() ? SYCLIntegrationHeader::kind_pointer : SYCLIntegrationHeader::kind_std_layout); @@ -3795,7 +3814,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { // The compiler generated kernel argument used to initialize SYCL 2020 // specialization constants, `specialization_constants_buffer`, should // have corresponding entry in integration header. - ASTContext &Context = SemaRef.getASTContext(); + ASTContext &Context = SemaSYCLRef.getASTContext(); // Offset is zero since kernel_handler argument is not part of // kernel object (i.e. it is not captured) addParam(Context.getPointerType(Context.CharTy), @@ -3832,7 +3851,7 @@ class SyclKernelIntHeaderCreator : public SyclKernelFieldHandler { } bool nextElement(QualType ET, uint64_t Index) final { - int64_t Size = SemaRef.getASTContext().getTypeSizeInChars(ET).getQuantity(); + int64_t Size = SemaSYCLRef.getASTContext().getTypeSizeInChars(ET).getQuantity(); CurOffset = ArrayBaseOffsets.back() + Size * Index; return true; } @@ -3851,7 +3870,7 @@ class SyclKernelIntFooterCreator : public SyclKernelFieldHandler { SYCLIntegrationFooter &Footer; public: - SyclKernelIntFooterCreator(Sema &S, SYCLIntegrationFooter &F) + SyclKernelIntFooterCreator(SemaSYCL &S, SYCLIntegrationFooter &F) : SyclKernelFieldHandler(S), Footer(F) { (void)Footer; // workaround for unused field warning } @@ -3862,7 +3881,7 @@ class SyclKernelIntFooterCreator : public SyclKernelFieldHandler { class SYCLKernelNameTypeVisitor : public TypeVisitor, public ConstTemplateArgumentVisitor { - Sema &S; + SemaSYCL &S; SourceLocation KernelInvocationFuncLoc; QualType KernelNameType; using InnerTypeVisitor = TypeVisitor; @@ -3877,7 +3896,7 @@ class SYCLKernelNameTypeVisitor } public: - SYCLKernelNameTypeVisitor(Sema &S, SourceLocation KernelInvocationFuncLoc, + SYCLKernelNameTypeVisitor(SemaSYCL &S, SourceLocation KernelInvocationFuncLoc, QualType KernelNameType, bool IsUnnamedKernel) : S(S), KernelInvocationFuncLoc(KernelInvocationFuncLoc), KernelNameType(KernelNameType), IsUnnamedKernel(IsUnnamedKernel) {} @@ -4044,7 +4063,7 @@ class SYCLKernelNameTypeVisitor } }; -void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, +void SemaSYCL::CheckSYCLKernelCall(FunctionDecl *KernelFunc, ArrayRef Args) { QualType KernelNameType = calculateKernelNameType(getASTContext(), KernelFunc); @@ -4055,7 +4074,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, // FIXME: In place until the library works around its 'host' invocation // issues. - if (!LangOpts.SYCLIsDevice) + if (SemaRef.LangOpts.SYCLIsDevice) return; const CXXRecordDecl *KernelObj = @@ -4079,11 +4098,11 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, QualType KernelParamTy = KernelFunc->getParamDecl(0)->getType(); if (KernelParamTy->isReferenceType()) { // passing by reference, so emit warning if not using SYCL 2020 - if (LangOpts.getSYCLVersion() < LangOptions::SYCL_2020) + if (SemaRef.LangOpts.getSYCLVersion() < LangOptions::SYCL_2020) Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_reference_future); } else { // passing by value. emit warning if using SYCL 2020 or greater - if (LangOpts.getSYCLVersion() > LangOptions::SYCL_2017) + if (SemaRef.LangOpts.getSYCLVersion() > LangOptions::SYCL_2017) Diag(KernelFunc->getLocation(), diag::warn_sycl_pass_by_value_deprecated); } @@ -4113,7 +4132,7 @@ void Sema::CheckSYCLKernelCall(FunctionDecl *KernelFunc, // For a wrapped parallel_for, copy attributes from original // kernel to wrapped kernel. -void Sema::copySYCLKernelAttrs(CXXMethodDecl *CallOperator) { +void SemaSYCL::copySYCLKernelAttrs(CXXMethodDecl *CallOperator) { // Get the operator() function of the wrapper. assert(CallOperator && "invalid kernel object"); @@ -4161,7 +4180,7 @@ void Sema::copySYCLKernelAttrs(CXXMethodDecl *CallOperator) { } } -void Sema::SetSYCLKernelNames() { +void SemaSYCL::SetSYCLKernelNames() { std::unique_ptr MangleCtx( getASTContext().createMangleContext()); // We assume the list of KernelDescs is the complete list of kernels needing @@ -4178,9 +4197,10 @@ void Sema::SetSYCLKernelNames() { StableName); // Set name of generated kernel. - Pair.second->setDeclName(&Context.Idents.get(KernelName)); + Pair.second->setDeclName(&getASTContext().Idents.get(KernelName)); // Update the AsmLabel for this generated kernel. - Pair.second->addAttr(AsmLabelAttr::CreateImplicit(Context, KernelName)); + Pair.second->addAttr( + AsmLabelAttr::CreateImplicit(getASTContext(), KernelName)); } } @@ -4206,7 +4226,7 @@ void Sema::SetSYCLKernelNames() { // } // // -void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, +void SemaSYCL::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, MangleContext &MC) { // The first argument to the KernelCallerFunc is the lambda object. const CXXRecordDecl *KernelObj = @@ -4254,7 +4274,7 @@ void Sema::ConstructOpenCLKernel(FunctionDecl *KernelCallerFunc, CallOperator); SyclKernelIntHeaderCreator int_header( IsSIMDKernel, *this, getSyclIntegrationHeader(), KernelObj, - calculateKernelNameType(Context, KernelCallerFunc), KernelCallerFunc); + calculateKernelNameType(getASTContext(), KernelCallerFunc), KernelCallerFunc); SyclKernelIntFooterCreator int_footer(*this, getSyclIntegrationFooter()); SyclOptReportCreator opt_report(*this, kernel_decl, KernelObj->getLocation()); @@ -4317,11 +4337,11 @@ static SourceLocation GetSubGroupLoc(const FunctionDecl *FD) { return SourceLocation{}; } -static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel, +static void CheckSYCL2020SubGroupSizes(SemaSYCL &S, FunctionDecl *SYCLKernel, const FunctionDecl *FD) { // If they are the same, no error. - if (CalcEffectiveSubGroup(S.Context, S.getLangOpts(), SYCLKernel) == - CalcEffectiveSubGroup(S.Context, S.getLangOpts(), FD)) + if (CalcEffectiveSubGroup(S.getASTContext(), S.getLangOpts(), SYCLKernel) == + CalcEffectiveSubGroup(S.getASTContext(), S.getLangOpts(), FD)) return; // No need to validate __spirv routines here since they @@ -4370,7 +4390,7 @@ static void CheckSYCL2020SubGroupSizes(Sema &S, FunctionDecl *SYCLKernel, // self-documentation purposes that it would be nice to be able to repeat these // on subsequent functions. static void CheckSYCL2020Attributes( - Sema &S, FunctionDecl *SYCLKernel, FunctionDecl *KernelBody, + SemaSYCL &S, FunctionDecl *SYCLKernel, FunctionDecl *KernelBody, const llvm::SmallPtrSetImpl &CalledFuncs) { if (KernelBody) { @@ -4416,7 +4436,7 @@ static void CheckSYCL2020Attributes( } static void PropagateAndDiagnoseDeviceAttr( - Sema &S, const SingleDeviceFunctionTracker &Tracker, Attr *A, + SemaSYCL &S, const SingleDeviceFunctionTracker &Tracker, Attr *A, FunctionDecl *SYCLKernel, FunctionDecl *KernelBody) { switch (A->getKind()) { case attr::Kind::IntelReqdSubGroupSize: { @@ -4446,9 +4466,9 @@ static void PropagateAndDiagnoseDeviceAttr( case attr::Kind::SYCLReqdWorkGroupSize: { auto *RWGSA = cast(A); if (auto *Existing = SYCLKernel->getAttr()) { - if (S.AnyWorkGroupSizesDiffer(Existing->getXDim(), Existing->getYDim(), - Existing->getZDim(), RWGSA->getXDim(), - RWGSA->getYDim(), RWGSA->getZDim())) { + if (S.SemaRef.AnyWorkGroupSizesDiffer( + Existing->getXDim(), Existing->getYDim(), Existing->getZDim(), + RWGSA->getXDim(), RWGSA->getYDim(), RWGSA->getZDim())) { S.Diag(SYCLKernel->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(Existing->getLocation(), diag::note_conflicting_attribute); @@ -4457,7 +4477,7 @@ static void PropagateAndDiagnoseDeviceAttr( } } else if (auto *Existing = SYCLKernel->getAttr()) { - if (S.CheckMaxAllowedWorkGroupSize( + if (S.SemaRef.CheckMaxAllowedWorkGroupSize( RWGSA->getXDim(), RWGSA->getYDim(), RWGSA->getZDim(), Existing->getXDim(), Existing->getYDim(), Existing->getZDim())) { S.Diag(SYCLKernel->getLocation(), @@ -4476,9 +4496,9 @@ static void PropagateAndDiagnoseDeviceAttr( case attr::Kind::SYCLWorkGroupSizeHint: { auto *WGSH = cast(A); if (auto *Existing = SYCLKernel->getAttr()) { - if (S.AnyWorkGroupSizesDiffer(Existing->getXDim(), Existing->getYDim(), - Existing->getZDim(), WGSH->getXDim(), - WGSH->getYDim(), WGSH->getZDim())) { + if (S.SemaRef.AnyWorkGroupSizesDiffer( + Existing->getXDim(), Existing->getYDim(), Existing->getZDim(), + WGSH->getXDim(), WGSH->getYDim(), WGSH->getZDim())) { S.Diag(SYCLKernel->getLocation(), diag::err_conflicting_sycl_kernel_attributes); S.Diag(Existing->getLocation(), diag::note_conflicting_attribute); @@ -4492,7 +4512,7 @@ static void PropagateAndDiagnoseDeviceAttr( case attr::Kind::SYCLIntelMaxWorkGroupSize: { auto *SIMWGSA = cast(A); if (auto *Existing = SYCLKernel->getAttr()) { - if (S.CheckMaxAllowedWorkGroupSize( + if (S.SemaRef.CheckMaxAllowedWorkGroupSize( Existing->getXDim(), Existing->getYDim(), Existing->getZDim(), SIMWGSA->getXDim(), SIMWGSA->getYDim(), SIMWGSA->getZDim())) { S.Diag(SYCLKernel->getLocation(), @@ -4547,7 +4567,7 @@ static void PropagateAndDiagnoseDeviceAttr( } } -void Sema::MarkDevices() { +void SemaSYCL::MarkDevices() { // This Tracker object ensures that the SyclDeviceDecls collection includes // the SYCL_EXTERNAL functions, and manages the diagnostics for all of the // functions in the kernel. @@ -4566,7 +4586,7 @@ void Sema::MarkDevices() { for (auto *A : T.GetCollectedAttributes()) PropagateAndDiagnoseDeviceAttr(*this, T, A, T.GetSYCLKernel(), T.GetKernelBody()); - CheckSYCLAddIRAttributesFunctionAttrConflicts(T.GetSYCLKernel()); + SemaRef.CheckSYCLAddIRAttributesFunctionAttrConflicts(T.GetSYCLKernel()); } } @@ -4575,19 +4595,19 @@ void Sema::MarkDevices() { // ----------------------------------------------------------------------------- Sema::SemaDiagnosticBuilder -Sema::SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID, +SemaSYCL::DiagIfDeviceCode(SourceLocation Loc, unsigned DiagID, DeviceDiagnosticReason Reason) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); - FunctionDecl *FD = dyn_cast(getCurLexicalContext()); + FunctionDecl *FD = dyn_cast(SemaRef.getCurLexicalContext()); SemaDiagnosticBuilder::Kind DiagKind = [this, FD, Reason] { if (DiagnosingSYCLKernel) return SemaDiagnosticBuilder::K_ImmediateWithCallStack; if (!FD) return SemaDiagnosticBuilder::K_Nop; - if (getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted) { + if (SemaRef.getEmissionStatus(FD) == Sema::FunctionEmissionStatus::Emitted) { // Skip the diagnostic if we know it won't be emitted. - if ((getEmissionReason(FD) & Reason) == + if ((SemaRef.getEmissionReason(FD) & Reason) == Sema::DeviceDiagnosticReason::None) return SemaDiagnosticBuilder::K_Nop; @@ -4595,10 +4615,10 @@ Sema::SYCLDiagIfDeviceCode(SourceLocation Loc, unsigned DiagID, } return SemaDiagnosticBuilder::K_Deferred; }(); - return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, *this, Reason); + return SemaDiagnosticBuilder(DiagKind, Loc, DiagID, FD, SemaRef, Reason); } -void Sema::deepTypeCheckForSYCLDevice(SourceLocation UsedAt, +void SemaSYCL::deepTypeCheckForDevice(SourceLocation UsedAt, llvm::DenseSet Visited, ValueDecl *DeclToCheck) { assert(getLangOpts().SYCLIsDevice && @@ -4610,18 +4630,18 @@ void Sema::deepTypeCheckForSYCLDevice(SourceLocation UsedAt, auto Check = [&](QualType TypeToCheck, const ValueDecl *D) { bool ErrorFound = false; if (isZeroSizedArray(*this, TypeToCheck)) { - SYCLDiagIfDeviceCode(UsedAt, diag::err_typecheck_zero_array_size) << 1; + DiagIfDeviceCode(UsedAt, diag::err_typecheck_zero_array_size) << 1; ErrorFound = true; } // Checks for other types can also be done here. if (ErrorFound) { if (NeedToEmitNotes) { if (auto *FD = dyn_cast(D)) - SYCLDiagIfDeviceCode(FD->getLocation(), - diag::note_illegal_field_declared_here) + DiagIfDeviceCode(FD->getLocation(), + diag::note_illegal_field_declared_here) << FD->getType()->isPointerType() << FD->getType(); else - SYCLDiagIfDeviceCode(D->getLocation(), diag::note_declared_at); + DiagIfDeviceCode(D->getLocation(), diag::note_declared_at); } } @@ -4652,8 +4672,8 @@ void Sema::deepTypeCheckForSYCLDevice(SourceLocation UsedAt, auto EmitHistory = [&]() { // The first element is always nullptr. for (uint64_t Index = 1; Index < History.size(); ++Index) { - SYCLDiagIfDeviceCode(History[Index]->getLocation(), - diag::note_within_field_of_type) + DiagIfDeviceCode(History[Index]->getLocation(), + diag::note_within_field_of_type) << History[Index]->getType(); } }; @@ -4690,7 +4710,7 @@ void Sema::deepTypeCheckForSYCLDevice(SourceLocation UsedAt, } while (!StackForRecursion.empty()); } -void Sema::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, +void SemaSYCL::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, const FunctionDecl *Callee, SourceLocation Loc, DeviceDiagnosticReason Reason) { @@ -4716,14 +4736,14 @@ void Sema::finalizeSYCLDelayedAnalysis(const FunctionDecl *Caller, // this undefined function is used to trigger a compiling error. if (!Callee->isDefined() && !Callee->getBuiltinID() && !Callee->isReplaceableGlobalAllocationFunction() && - !isSYCLUndefinedAllowed(Callee, getSourceManager())) { - Diag(Loc, diag::err_sycl_restrict) << Sema::KernelCallUndefinedFunction; + !isSYCLUndefinedAllowed(Callee, SemaRef.getSourceManager())) { + Diag(Loc, diag::err_sycl_restrict) << SemaSYCL::KernelCallUndefinedFunction; Diag(Callee->getLocation(), diag::note_previous_decl) << Callee; Diag(Caller->getLocation(), diag::note_called_by) << Caller; } } -bool Sema::checkAllowedSYCLInitializer(VarDecl *VD) { +bool SemaSYCL::checkAllowedSYCLInitializer(VarDecl *VD) { assert(getLangOpts().SYCLIsDevice && "Should only be called during SYCL compilation"); @@ -4732,8 +4752,8 @@ bool Sema::checkAllowedSYCLInitializer(VarDecl *VD) { const Expr *Init = VD->getInit(); bool ValueDependent = Init && Init->isValueDependent(); - bool isConstantInit = - Init && !ValueDependent && Init->isConstantInitializer(Context, false); + bool isConstantInit = Init && !ValueDependent && + Init->isConstantInitializer(getASTContext(), false); if (!VD->isConstexpr() && Init && !ValueDependent && !isConstantInit) return false; @@ -5317,8 +5337,9 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { for (const KernelDesc &K : KernelDescs) { const size_t N = K.Params.size(); - PresumedLoc PLoc = S.Context.getSourceManager().getPresumedLoc( - S.Context.getSourceManager() + PresumedLoc PLoc = S.getASTContext().getSourceManager().getPresumedLoc( + S.getASTContext() + .getSourceManager() .getExpansionRange(K.KernelLocation) .getEnd()); if (K.IsUnnamedKernel) { @@ -5383,10 +5404,10 @@ void SYCLIntegrationHeader::emit(raw_ostream &O) { O << " return 0;\n"; O << "#endif\n"; O << " }\n"; - StringRef ReturnType = - (S.Context.getTargetInfo().getInt64Type() == TargetInfo::SignedLong) - ? "long" - : "long long"; + StringRef ReturnType = (S.getASTContext().getTargetInfo().getInt64Type() == + TargetInfo::SignedLong) + ? "long" + : "long long"; O << " // Returns the size of the kernel object in bytes.\n"; O << " __SYCL_DLL_LOCAL\n"; O << " static constexpr " << ReturnType << " getKernelSize() { return " @@ -5445,7 +5466,7 @@ void SYCLIntegrationHeader::addSpecConstant(StringRef IDName, QualType IDType) { SpecConsts.emplace_back(std::make_pair(IDType, IDName.str())); } -SYCLIntegrationHeader::SYCLIntegrationHeader(Sema &S) : S(S) {} +SYCLIntegrationHeader::SYCLIntegrationHeader(SemaSYCL &S) : S(S) {} void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) { // Variable template declaration can result in an error case of 'nullptr' @@ -5462,8 +5483,8 @@ void SYCLIntegrationFooter::addVarDecl(const VarDecl *VD) { if (isa(VD)) return; // Step 1: ensure that this is of the correct type template specialization. - if (!Sema::isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) && - !Sema::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && + if (!SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) && + !SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && !S.isTypeDecoratedWithDeclAttribute( VD->getType())) { // Handle the case where this could be a deduced type, such as a deduction @@ -5659,8 +5680,8 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { // Skip if this isn't a SpecIdType, DeviceGlobal, or HostPipe. This // can happen if it was a deduced type. - if (!Sema::isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) && - !Sema::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && + if (!SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::specialization_id) && + !SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe) && !S.isTypeDecoratedWithDeclAttribute( VD->getType())) continue; @@ -5691,7 +5712,7 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { DeviceGlobOS << SYCLUniqueStableIdExpr::ComputeName(S.getASTContext(), VD); DeviceGlobOS << "\");\n"; - } else if (Sema::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe)) { + } else if (SemaSYCL::isSyclType(VD->getType(), SYCLTypeAttr::host_pipe)) { HostPipesEmitted = true; HostPipesOS << "host_pipe_map::add("; HostPipesOS << "(void *)&"; @@ -5766,3 +5787,64 @@ bool SYCLIntegrationFooter::emit(raw_ostream &OS) { return true; } + +ExprResult SemaSYCL::BuildUniqueStableIdExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, Expr *E) { + if (!E->isInstantiationDependent()) { + // Special handling to get us better error messages for a member variable. + if (auto *ME = dyn_cast(E->IgnoreUnlessSpelledInSource())) { + if (isa(ME->getMemberDecl())) + Diag(E->getExprLoc(), diag::err_unique_stable_id_global_storage); + else + Diag(E->getExprLoc(), diag::err_unique_stable_id_expected_var); + return ExprError(); + } + + auto *DRE = dyn_cast(E->IgnoreUnlessSpelledInSource()); + + if (!DRE || !isa_and_nonnull(DRE->getDecl())) { + Diag(E->getExprLoc(), diag::err_unique_stable_id_expected_var); + return ExprError(); + } + + auto *Var = cast(DRE->getDecl()); + + if (!Var->hasGlobalStorage()) { + Diag(E->getExprLoc(), diag::err_unique_stable_id_global_storage); + return ExprError(); + } + } + + return SYCLUniqueStableIdExpr::Create(getASTContext(), OpLoc, LParen, RParen, + E); +} + +ExprResult SemaSYCL::ActOnUniqueStableIdExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, Expr *E) { + return BuildUniqueStableIdExpr(OpLoc, LParen, RParen, E); +} + +ExprResult SemaSYCL::BuildUniqueStableNameExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, + TypeSourceInfo *TSI) { + return SYCLUniqueStableNameExpr::Create(getASTContext(), OpLoc, LParen, + RParen, TSI); +} + +ExprResult SemaSYCL::ActOnUniqueStableNameExpr(SourceLocation OpLoc, + SourceLocation LParen, + SourceLocation RParen, + ParsedType ParsedTy) { + TypeSourceInfo *TSI = nullptr; + QualType Ty = SemaRef.GetTypeFromParser(ParsedTy, &TSI); + + if (Ty.isNull()) + return ExprError(); + if (!TSI) + TSI = getASTContext().getTrivialTypeSourceInfo(Ty, LParen); + + return BuildUniqueStableNameExpr(OpLoc, LParen, RParen, TSI); +} diff --git a/clang/lib/Sema/SemaStmt.cpp b/clang/lib/Sema/SemaStmt.cpp index a9711a7397bc7..bb039c876c9ca 100644 --- a/clang/lib/Sema/SemaStmt.cpp +++ b/clang/lib/Sema/SemaStmt.cpp @@ -34,6 +34,7 @@ #include "clang/Sema/Scope.h" #include "clang/Sema/ScopeInfo.h" #include "clang/Sema/SemaInternal.h" +#include "clang/Sema/SemaSYCL.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/ADT/DenseMap.h" #include "llvm/ADT/STLExtras.h" @@ -4580,8 +4581,8 @@ StmtResult Sema::ActOnCXXTryBlock(SourceLocation TryLoc, Stmt *TryBlock, // Exceptions aren't allowed in SYCL device code. if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(TryLoc, diag::err_sycl_restrict) - << KernelUseExceptions; + SYCL().DiagIfDeviceCode(TryLoc, diag::err_sycl_restrict) + << SemaSYCL::KernelUseExceptions; if (getCurScope() && getCurScope()->isOpenMPSimdDirectiveScope()) Diag(TryLoc, diag::err_omp_simd_region_cannot_use_stmt) << "try"; @@ -4693,8 +4694,8 @@ StmtResult Sema::ActOnSEHTryBlock(bool IsCXXTry, SourceLocation TryLoc, // Exceptions aren't allowed in SYCL device code. if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(TryLoc, diag::err_sycl_restrict) - << KernelUseExceptions; + SYCL().DiagIfDeviceCode(TryLoc, diag::err_sycl_restrict) + << SemaSYCL::KernelUseExceptions; FSI->setHasSEHTry(TryLoc); @@ -4712,7 +4713,7 @@ StmtResult Sema::ActOnSEHTryBlock(bool IsCXXTry, SourceLocation TryLoc, // Reject __try on unsupported targets. if (!Context.getTargetInfo().isSEHTrySupported()) { if (getLangOpts().SYCLIsDevice) - SYCLDiagIfDeviceCode(TryLoc, diag::err_seh_try_unsupported); + SYCL().DiagIfDeviceCode(TryLoc, diag::err_seh_try_unsupported); else Diag(TryLoc, diag::err_seh_try_unsupported); } diff --git a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp index 7d65092e6c920..fe11ee7b44839 100644 --- a/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp +++ b/clang/lib/Sema/SemaTemplateInstantiateDecl.cpp @@ -1732,7 +1732,7 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D, // adding the VarTemplateSpecializationDecl later. if (!InstantiatingVarTemplate) { if (SemaRef.getLangOpts().SYCLIsDevice && - SemaRef.isTypeDecoratedWithDeclAttribute( + SemaRef.SYCL().isTypeDecoratedWithDeclAttribute( Var->getType())) { if (!Var->hasGlobalStorage()) SemaRef.Diag(D->getLocation(), @@ -1756,13 +1756,13 @@ Decl *TemplateDeclInstantiator::VisitVarDecl(VarDecl *D, } } if (const auto *SYCLDevice = Var->getAttr()) { - if (!SemaRef.isTypeDecoratedWithDeclAttribute( + if (!SemaRef.SYCL().isTypeDecoratedWithDeclAttribute( Var->getType())) SemaRef.Diag(SYCLDevice->getLoc(), diag::err_sycl_attribute_not_device_global) << SYCLDevice; } - SemaRef.addSyclVarDecl(Var); + SemaRef.SYCL().addSyclVarDecl(Var); } if (Var->getTLSKind()) @@ -1859,7 +1859,7 @@ Decl *TemplateDeclInstantiator::VisitFieldDecl(FieldDecl *D) { // Static members are not processed here, so error out if we have a device // global without checking access modifier. if (SemaRef.getLangOpts().SYCLIsDevice) { - if (SemaRef.isTypeDecoratedWithDeclAttribute( + if (SemaRef.SYCL().isTypeDecoratedWithDeclAttribute( Field->getType())) { SemaRef.Diag(D->getLocation(), diag::err_sycl_device_global_incorrect_scope); @@ -5611,14 +5611,6 @@ void Sema::InstantiateFunctionDefinition(SourceLocation PointOfInstantiation, EnterExpressionEvaluationContext EvalContext( *this, Sema::ExpressionEvaluationContext::PotentiallyEvaluated); - Qualifiers ThisTypeQuals; - CXXRecordDecl *ThisContext = nullptr; - if (CXXMethodDecl *Method = dyn_cast(Function)) { - ThisContext = Method->getParent(); - ThisTypeQuals = Method->getMethodQualifiers(); - } - CXXThisScopeRAII ThisScope(*this, ThisContext, ThisTypeQuals); - // Introduce a new scope where local variable instantiations will be // recorded, unless we're actually a member function within a local // class, in which case we need to merge our results with the parent @@ -5821,7 +5813,7 @@ VarTemplateSpecializationDecl *Sema::BuildVarTemplateInstantiation( Instantiator.VisitVarTemplateSpecializationDecl( VarTemplate, FromVar, TemplateArgsInfo, Converted)); - addSyclVarDecl(VD); + SYCL().addSyclVarDecl(VD); return VD; } @@ -7010,7 +7002,7 @@ static void processFunctionInstantiation(Sema &S, if (!FD->isDefined()) return; if (S.LangOpts.SYCLIsDevice && FD->hasAttr()) - S.ConstructOpenCLKernel(FD, MC); + S.SYCL().ConstructOpenCLKernel(FD, MC); FD->setInstantiationIsPending(false); } diff --git a/clang/lib/Sema/SemaTemplateVariadic.cpp b/clang/lib/Sema/SemaTemplateVariadic.cpp index 903fbfd18e779..4909414c0c78d 100644 --- a/clang/lib/Sema/SemaTemplateVariadic.cpp +++ b/clang/lib/Sema/SemaTemplateVariadic.cpp @@ -1243,6 +1243,17 @@ std::optional Sema::getFullyPackExpandedSize(TemplateArgument Arg) { // expanded this pack expansion into the enclosing pack if we could. if (Elem.isPackExpansion()) return std::nullopt; + // Don't guess the size of unexpanded packs. The pack within a template + // argument may have yet to be of a PackExpansion type before we see the + // ellipsis in the annotation stage. + // + // This doesn't mean we would invalidate the optimization: Arg can be an + // unexpanded pack regardless of Elem's dependence. For instance, + // A TemplateArgument that contains either a SubstTemplateTypeParmPackType + // or SubstNonTypeTemplateParmPackExpr is always considered Unexpanded, but + // the underlying TemplateArgument thereof may not. + if (Elem.containsUnexpandedParameterPack()) + return std::nullopt; } return Pack.pack_size(); } diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h index 935903ba4e3c6..b976e5d41b95e 100644 --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -40,6 +40,7 @@ #include "clang/Sema/SemaDiagnostic.h" #include "clang/Sema/SemaInternal.h" #include "clang/Sema/SemaOpenACC.h" +#include "clang/Sema/SemaSYCL.h" #include "llvm/ADT/ArrayRef.h" #include "llvm/Support/ErrorHandling.h" #include @@ -2634,13 +2635,15 @@ class TreeTransform { SourceLocation LParen, SourceLocation RParen, TypeSourceInfo *TSI) { - return getSema().BuildSYCLUniqueStableNameExpr(OpLoc, LParen, RParen, TSI); + return getSema().SYCL().BuildUniqueStableNameExpr(OpLoc, LParen, RParen, + TSI); } ExprResult RebuildSYCLUniqueStableIdExpr(SourceLocation OpLoc, SourceLocation LParen, SourceLocation RParen, Expr *E) { - return getSema().BuildSYCLUniqueStableIdExpr(OpLoc, LParen, RParen, E); + return getSema().SYCL().BuildUniqueStableIdExpr(OpLoc, LParen, RParen, + E); } /// Build a new predefined expression. @@ -3271,22 +3274,22 @@ class TreeTransform { ExprResult RebuildSYCLBuiltinNumFieldsExpr(SourceLocation Loc, QualType SourceTy) { - return getSema().BuildSYCLBuiltinNumFieldsExpr(Loc, SourceTy); + return getSema().SYCL().BuildSYCLBuiltinNumFieldsExpr(Loc, SourceTy); } ExprResult RebuildSYCLBuiltinFieldTypeExpr(SourceLocation Loc, QualType SourceTy, Expr *Idx) { - return getSema().BuildSYCLBuiltinFieldTypeExpr(Loc, SourceTy, Idx); + return getSema().SYCL().BuildSYCLBuiltinFieldTypeExpr(Loc, SourceTy, Idx); } ExprResult RebuildSYCLBuiltinNumBasesExpr(SourceLocation Loc, QualType SourceTy) { - return getSema().BuildSYCLBuiltinNumBasesExpr(Loc, SourceTy); + return getSema().SYCL().BuildSYCLBuiltinNumBasesExpr(Loc, SourceTy); } ExprResult RebuildSYCLBuiltinBaseTypeExpr(SourceLocation Loc, QualType SourceTy, Expr *Idx) { - return getSema().BuildSYCLBuiltinBaseTypeExpr(Loc, SourceTy, Idx); + return getSema().SYCL().BuildSYCLBuiltinBaseTypeExpr(Loc, SourceTy, Idx); } /// Build a new C++ typeid(type) expression. @@ -3335,13 +3338,12 @@ class TreeTransform { /// Build a new C++ "this" expression. /// - /// By default, performs semantic analysis to build a new "this" expression. - /// Subclasses may override this routine to provide different behavior. + /// By default, builds a new "this" expression without performing any + /// semantic analysis. Subclasses may override this routine to provide + /// different behavior. ExprResult RebuildCXXThisExpr(SourceLocation ThisLoc, QualType ThisType, bool isImplicit) { - if (getSema().CheckCXXThisType(ThisLoc, ThisType)) - return ExprError(); return getSema().BuildCXXThisExpr(ThisLoc, ThisType, isImplicit); } diff --git a/clang/test/Driver/linker-wrapper-libs.c b/clang/test/Driver/linker-wrapper-libs.c index 9a78200d7d3cf..119e306857187 100644 --- a/clang/test/Driver/linker-wrapper-libs.c +++ b/clang/test/Driver/linker-wrapper-libs.c @@ -44,6 +44,8 @@ int bar() { return weak; } // RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.o -fembed-offload-object=%t.out // RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \ // RUN: --linker-path=/usr/bin/ld %t.o %t.a -o a.out 2>&1 \ +// RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \ +// RUN: --linker-path=/usr/bin/ld %t.a %t.o -o a.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=LIBRARY-RESOLVES // LIBRARY-RESOLVES: clang{{.*}} -o {{.*}}.img --target=amdgcn-amd-amdhsa -mcpu=gfx1030 {{.*}}.o {{.*}}.o @@ -66,6 +68,8 @@ int bar() { return weak; } // RUN: %clang -cc1 %s -triple x86_64-unknown-linux-gnu -emit-obj -o %t.o -fembed-offload-object=%t.out // RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \ // RUN: --linker-path=/usr/bin/ld %t.o %t.a -o a.out 2>&1 \ +// RUN: clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu --dry-run \ +// RUN: --linker-path=/usr/bin/ld %t.a %t.o -o a.out 2>&1 \ // RUN: | FileCheck %s --check-prefix=LIBRARY-GLOBAL // LIBRARY-GLOBAL: clang{{.*}} -o {{.*}}.img --target=amdgcn-amd-amdhsa -mcpu=gfx1030 {{.*}}.o {{.*}}.o diff --git a/clang/test/SemaTemplate/alias-templates.cpp b/clang/test/SemaTemplate/alias-templates.cpp index 8d7cc6118610a..ab5cad72faf1b 100644 --- a/clang/test/SemaTemplate/alias-templates.cpp +++ b/clang/test/SemaTemplate/alias-templates.cpp @@ -236,6 +236,29 @@ namespace PR14858 { void test_q(int (&a)[5]) { Q().f(&a); } } +namespace PR84220 { + +template class list {}; + +template struct foo_impl { + template using f = int; +}; + +template +using foo = typename foo_impl::template f; + +// We call getFullyPackExpandedSize at the annotation stage +// before parsing the ellipsis next to the foo. This happens before +// a PackExpansionType is formed for foo. +// getFullyPackExpandedSize shouldn't determine the value here. Otherwise, +// foo_impl would lose its dependency despite the template +// arguments being unsubstituted. +template using test = list...>; + +test a; + +} + namespace redecl { template using A = int; template using A = int; diff --git a/clang/test/SemaTemplate/ms-function-specialization-class-scope.cpp b/clang/test/SemaTemplate/ms-function-specialization-class-scope.cpp index 6977623a0816e..dcab9bfaeabcb 100644 --- a/clang/test/SemaTemplate/ms-function-specialization-class-scope.cpp +++ b/clang/test/SemaTemplate/ms-function-specialization-class-scope.cpp @@ -1,6 +1,7 @@ -// RUN: %clang_cc1 -fms-extensions -fsyntax-only -Wno-unused-value -verify %s -// RUN: %clang_cc1 -fms-extensions -fdelayed-template-parsing -fsyntax-only -Wno-unused-value -verify %s +// RUN: %clang_cc1 -fms-extensions -fsyntax-only -verify %s +// RUN: %clang_cc1 -fms-extensions -fdelayed-template-parsing -fsyntax-only -verify %s +// expected-no-diagnostics class A { public: template A(U p) {} @@ -75,42 +76,3 @@ struct S { int f<0>(int); }; } - -namespace UsesThis { - template - struct A { - int x; - - template - static void f(); - - template<> - void f() { - this->x; // expected-error {{invalid use of 'this' outside of a non-static member function}} - x; // expected-error {{invalid use of member 'x' in static member function}} - A::x; // expected-error {{invalid use of member 'x' in static member function}} - +x; // expected-error {{invalid use of member 'x' in static member function}} - +A::x; // expected-error {{invalid use of member 'x' in static member function}} - } - - template - void g(); - - template<> - void g() { - this->x; - x; - A::x; - +x; - +A::x; - } - - template - static auto h() -> A*; - - template<> - auto h() -> decltype(this); // expected-error {{'this' cannot be used in a static member function declaration}} - }; - - template struct A; // expected-note 2{{in instantiation of}} -} diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index ec0ed432c1335..c01effcc4bf24 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2063,9 +2063,9 @@ getDeviceInput(const ArgList &Args) { StringSaver Saver(Alloc); // Try to extract device code from the linker input files. - DenseMap> InputFiles; - DenseMap> Syms; bool WholeArchive = Args.hasArg(OPT_wholearchive_flag) ? true : false; + SmallVector ObjectFilesToExtract; + SmallVector ArchiveFilesToExtract; for (const opt::Arg *Arg : Args.filtered( OPT_INPUT, OPT_library, OPT_whole_archive, OPT_no_whole_archive)) { if (Arg->getOption().matches(OPT_whole_archive) || @@ -2113,50 +2113,87 @@ getDeviceInput(const ArgList &Args) { if (Error Err = extractOffloadBinaries(Buffer, Binaries)) return std::move(Err); - // We only extract archive members that are needed. - bool IsArchive = identify_magic(Buffer.getBuffer()) == file_magic::archive; - bool Extracted = true; - while (Extracted) { - Extracted = false; - for (OffloadFile &Binary : Binaries) { - // If the binary was previously extracted it will be set to null. - if (!Binary.getBinary()) + for (auto &OffloadFile : Binaries) { + if (identify_magic(Buffer.getBuffer()) == file_magic::archive && + !WholeArchive) + ArchiveFilesToExtract.emplace_back(std::move(OffloadFile)); + else + ObjectFilesToExtract.emplace_back(std::move(OffloadFile)); + } + } + + // Link all standard input files and update the list of symbols. + DenseMap> InputFiles; + DenseMap> Syms; + for (OffloadFile &Binary : ObjectFilesToExtract) { + if (!Binary.getBinary()) + continue; + + SmallVector CompatibleTargets = {Binary}; + for (const auto &[ID, Input] : InputFiles) + if (object::areTargetsCompatible(Binary, ID)) + CompatibleTargets.emplace_back(ID); + + for (const auto &[Index, ID] : llvm::enumerate(CompatibleTargets)) { + Expected ExtractOrErr = getSymbols( + Binary.getBinary()->getImage(), Binary.getBinary()->getOffloadKind(), + /*IsArchive=*/false, Saver, Syms[ID]); + if (!ExtractOrErr) + return ExtractOrErr.takeError(); + + // If another target needs this binary it must be copied instead. + if (Index == CompatibleTargets.size() - 1) + InputFiles[ID].emplace_back(std::move(Binary)); + else + InputFiles[ID].emplace_back(Binary.copy()); + } + } + + // Archive members only extract if they define needed symbols. We do this + // after every regular input file so that libraries may be included out of + // order. This follows 'ld.lld' semantics which are more lenient. + bool Extracted = true; + while (Extracted) { + Extracted = false; + for (OffloadFile &Binary : ArchiveFilesToExtract) { + // If the binary was previously extracted it will be set to null. + if (!Binary.getBinary()) + continue; + + SmallVector CompatibleTargets = {Binary}; + for (const auto &[ID, Input] : InputFiles) + if (object::areTargetsCompatible(Binary, ID)) + CompatibleTargets.emplace_back(ID); + + for (const auto &[Index, ID] : llvm::enumerate(CompatibleTargets)) { + // Only extract an if we have an an object matching this target. + if (!InputFiles.count(ID)) continue; - SmallVector CompatibleTargets = {Binary}; - for (const auto &[ID, Input] : InputFiles) - if (object::areTargetsCompatible(Binary, ID)) - CompatibleTargets.emplace_back(ID); - - for (const auto &[Index, ID] : llvm::enumerate(CompatibleTargets)) { - // Only extract an if we have an an object matching this target. - if (IsArchive && !WholeArchive && !InputFiles.count(ID)) - continue; - - Expected ExtractOrErr = getSymbols( - Binary.getBinary()->getImage(), - Binary.getBinary()->getOffloadKind(), IsArchive, Saver, Syms[ID]); - if (!ExtractOrErr) - return ExtractOrErr.takeError(); - - Extracted = !WholeArchive && *ExtractOrErr; - - // Skip including the file if it is an archive that does not resolve - // any symbols. - if (IsArchive && !WholeArchive && !Extracted) - continue; - - // If another target needs this binary it must be copied instead. - if (Index == CompatibleTargets.size() - 1) - InputFiles[ID].emplace_back(std::move(Binary)); - else - InputFiles[ID].emplace_back(Binary.copy()); - } + Expected ExtractOrErr = + getSymbols(Binary.getBinary()->getImage(), + Binary.getBinary()->getOffloadKind(), /*IsArchive=*/true, + Saver, Syms[ID]); + if (!ExtractOrErr) + return ExtractOrErr.takeError(); - // If we extracted any files we need to check all the symbols again. - if (Extracted) - break; + Extracted = *ExtractOrErr; + + // Skip including the file if it is an archive that does not resolve + // any symbols. + if (!Extracted) + continue; + + // If another target needs this binary it must be copied instead. + if (Index == CompatibleTargets.size() - 1) + InputFiles[ID].emplace_back(std::move(Binary)); + else + InputFiles[ID].emplace_back(Binary.copy()); } + + // If we extracted any files we need to check all the symbols again. + if (Extracted) + break; } } diff --git a/lld/COFF/ICF.cpp b/lld/COFF/ICF.cpp index 013ffcfb3d5d1..b899a25324239 100644 --- a/lld/COFF/ICF.cpp +++ b/lld/COFF/ICF.cpp @@ -178,7 +178,7 @@ bool ICF::equalsConstant(const SectionChunk *a, const SectionChunk *b) { a->getSectionName() == b->getSectionName() && a->header->SizeOfRawData == b->header->SizeOfRawData && a->checksum == b->checksum && a->getContents() == b->getContents() && - assocEquals(a, b); + a->getMachine() == b->getMachine() && assocEquals(a, b); } // Compare "moving" part of two sections, namely relocation targets. diff --git a/lld/test/COFF/arm64x-icf.s b/lld/test/COFF/arm64x-icf.s new file mode 100644 index 0000000000000..c8df21d3e4969 --- /dev/null +++ b/lld/test/COFF/arm64x-icf.s @@ -0,0 +1,37 @@ +// REQUIRES: aarch64 +// RUN: split-file %s %t.dir && cd %t.dir + +// RUN: llvm-mc -filetype=obj -triple=arm64ec-windows func-arm64ec.s -o func-arm64ec.obj +// RUN: llvm-mc -filetype=obj -triple=aarch64-windows func-arm64.s -o func-arm64.obj +// RUN: lld-link -machine:arm64x -dll -noentry -out:out.dll func-arm64ec.obj func-arm64.obj +// RUN: llvm-objdump -d out.dll | FileCheck %s + +// CHECK: 0000000180001000 <.text>: +// CHECK-NEXT: 180001000: 52800020 mov w0, #0x1 // =1 +// CHECK-NEXT: 180001004: d65f03c0 ret +// CHECK-NEXT: ... +// CHECK-NEXT: 180002000: 52800020 mov w0, #0x1 // =1 +// CHECK-NEXT: 180002004: d65f03c0 ret + + +#--- func-arm64.s + .section .text,"xr",discard,func + .globl func + .p2align 2 +func: + mov w0, #1 + ret + + .data + .rva func + +#--- func-arm64ec.s + .section .text,"xr",discard,"#func" + .globl "#func" + .p2align 2 +"#func": + mov w0, #1 + ret + + .data + .rva "#func" diff --git a/llvm/lib/Analysis/ScalarEvolution.cpp b/llvm/lib/Analysis/ScalarEvolution.cpp index e030b9fc7dac4..9fcce797f5597 100644 --- a/llvm/lib/Analysis/ScalarEvolution.cpp +++ b/llvm/lib/Analysis/ScalarEvolution.cpp @@ -928,11 +928,9 @@ static const SCEV *BinomialCoefficient(const SCEV *It, unsigned K, APInt OddFactorial(W, 1); unsigned T = 1; for (unsigned i = 3; i <= K; ++i) { - APInt Mult(W, i); - unsigned TwoFactors = Mult.countr_zero(); + unsigned TwoFactors = countr_zero(i); T += TwoFactors; - Mult.lshrInPlace(TwoFactors); - OddFactorial *= Mult; + OddFactorial *= (i >> TwoFactors); } // We need at least W + T bits for the multiplication step diff --git a/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp b/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp index c3dcf73b0b762..22ef9b5fb994e 100644 --- a/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp +++ b/llvm/lib/Transforms/Vectorize/SLPVectorizer.cpp @@ -2325,19 +2325,17 @@ class BoUpSLP { ~BoUpSLP(); private: - /// Determine if a vectorized value \p V in can be demoted to - /// a smaller type with a truncation. We collect the values that will be - /// demoted in ToDemote and additional roots that require investigating in - /// Roots. - /// \param DemotedConsts list of Instruction/OperandIndex pairs that are - /// constant and to be demoted. Required to correctly identify constant nodes - /// to be demoted. - bool collectValuesToDemote( - Value *V, bool IsProfitableToDemoteRoot, unsigned &BitWidth, - SmallVectorImpl &ToDemote, - DenseMap> &DemotedConsts, - DenseSet &Visited, unsigned &MaxDepthLevel, - bool &IsProfitableToDemote, bool IsTruncRoot) const; + /// Determine if a node \p E in can be demoted to a smaller type with a + /// truncation. We collect the entries that will be demoted in ToDemote. + /// \param E Node for analysis + /// \param ToDemote indices of the nodes to be demoted. + bool collectValuesToDemote(const TreeEntry &E, bool IsProfitableToDemoteRoot, + unsigned &BitWidth, + SmallVectorImpl &ToDemote, + DenseSet &Visited, + unsigned &MaxDepthLevel, + bool &IsProfitableToDemote, + bool IsTruncRoot) const; /// Check if the operands on the edges \p Edges of the \p UserTE allows /// reordering (i.e. the operands can be reordered because they have only one @@ -14126,20 +14124,17 @@ unsigned BoUpSLP::getVectorElementSize(Value *V) { return Width; } -// Determine if a value V in a vectorizable expression Expr can be demoted to a -// smaller type with a truncation. We collect the values that will be demoted -// in ToDemote and additional roots that require investigating in Roots. bool BoUpSLP::collectValuesToDemote( - Value *V, bool IsProfitableToDemoteRoot, unsigned &BitWidth, - SmallVectorImpl &ToDemote, - DenseMap> &DemotedConsts, - DenseSet &Visited, unsigned &MaxDepthLevel, - bool &IsProfitableToDemote, bool IsTruncRoot) const { + const TreeEntry &E, bool IsProfitableToDemoteRoot, unsigned &BitWidth, + SmallVectorImpl &ToDemote, DenseSet &Visited, + unsigned &MaxDepthLevel, bool &IsProfitableToDemote, + bool IsTruncRoot) const { // We can always demote constants. - if (isa(V)) + if (all_of(E.Scalars, IsaPred)) return true; - if (DL->getTypeSizeInBits(V->getType()) == BitWidth) { + unsigned OrigBitWidth = DL->getTypeSizeInBits(E.Scalars.front()->getType()); + if (OrigBitWidth == BitWidth) { MaxDepthLevel = 1; return true; } @@ -14150,7 +14145,6 @@ bool BoUpSLP::collectValuesToDemote( auto IsPotentiallyTruncated = [&](Value *V, unsigned &BitWidth) -> bool { if (MultiNodeScalars.contains(V)) return false; - uint32_t OrigBitWidth = DL->getTypeSizeInBits(V->getType()); if (OrigBitWidth > BitWidth) { APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); if (MaskedValueIsZero(V, Mask, SimplifyQuery(*DL))) @@ -14168,47 +14162,50 @@ bool BoUpSLP::collectValuesToDemote( BitWidth = std::max(BitWidth, BitWidth1); return BitWidth > 0 && OrigBitWidth >= (BitWidth * 2); }; - auto FinalAnalysis = [&](const TreeEntry *ITE = nullptr) { + using namespace std::placeholders; + auto FinalAnalysis = [&]() { if (!IsProfitableToDemote) return false; - return (ITE && ITE->UserTreeIndices.size() > 1) || - IsPotentiallyTruncated(V, BitWidth); + bool Res = all_of( + E.Scalars, std::bind(IsPotentiallyTruncated, _1, std::ref(BitWidth))); + // Gather demoted constant operands. + if (Res && E.State == TreeEntry::NeedToGather && + all_of(E.Scalars, IsaPred)) + ToDemote.push_back(E.Idx); + return Res; }; // TODO: improve handling of gathered values and others. - auto *I = dyn_cast(V); - const TreeEntry *ITE = I ? getTreeEntry(I) : nullptr; - if (!ITE || !Visited.insert(I).second || MultiNodeScalars.contains(I) || - all_of(I->users(), [&](User *U) { - return isa(U) && !getTreeEntry(U); + if (E.State == TreeEntry::NeedToGather || !Visited.insert(&E).second || + any_of(E.Scalars, [&](Value *V) { + return all_of(V->users(), [&](User *U) { + return isa(U) && !getTreeEntry(U); + }); })) return FinalAnalysis(); - if (!all_of(I->users(), - [=](User *U) { - return getTreeEntry(U) || - (UserIgnoreList && UserIgnoreList->contains(U)) || - (U->getType()->isSized() && - !U->getType()->isScalableTy() && - DL->getTypeSizeInBits(U->getType()) <= BitWidth); - }) && - !IsPotentiallyTruncated(I, BitWidth)) + if (any_of(E.Scalars, [&](Value *V) { + return !all_of(V->users(), [=](User *U) { + return getTreeEntry(U) || + (UserIgnoreList && UserIgnoreList->contains(U)) || + (U->getType()->isSized() && !U->getType()->isScalableTy() && + DL->getTypeSizeInBits(U->getType()) <= BitWidth); + }) && !IsPotentiallyTruncated(V, BitWidth); + })) return false; - unsigned Start = 0; - unsigned End = I->getNumOperands(); - - auto ProcessOperands = [&](ArrayRef Operands, bool &NeedToExit) { + auto ProcessOperands = [&](ArrayRef Operands, + bool &NeedToExit) { NeedToExit = false; unsigned InitLevel = MaxDepthLevel; - for (Value *IncValue : Operands) { + for (const TreeEntry *Op : Operands) { unsigned Level = InitLevel; - if (!collectValuesToDemote(IncValue, IsProfitableToDemoteRoot, BitWidth, - ToDemote, DemotedConsts, Visited, Level, - IsProfitableToDemote, IsTruncRoot)) { + if (!collectValuesToDemote(*Op, IsProfitableToDemoteRoot, BitWidth, + ToDemote, Visited, Level, IsProfitableToDemote, + IsTruncRoot)) { if (!IsProfitableToDemote) return false; NeedToExit = true; - if (!FinalAnalysis(ITE)) + if (!FinalAnalysis()) return false; continue; } @@ -14220,7 +14217,6 @@ bool BoUpSLP::collectValuesToDemote( [&](function_ref Checker, bool &NeedToExit) { // Try all bitwidth < OrigBitWidth. NeedToExit = false; - uint32_t OrigBitWidth = DL->getTypeSizeInBits(I->getType()); unsigned BestFailBitwidth = 0; for (; BitWidth < OrigBitWidth; BitWidth *= 2) { if (Checker(BitWidth, OrigBitWidth)) @@ -14241,18 +14237,20 @@ bool BoUpSLP::collectValuesToDemote( return false; }; auto TryProcessInstruction = - [&](Instruction *I, const TreeEntry &ITE, unsigned &BitWidth, - ArrayRef Operands = std::nullopt, + [&](unsigned &BitWidth, + ArrayRef Operands = std::nullopt, function_ref Checker = {}) { if (Operands.empty()) { if (!IsTruncRoot) MaxDepthLevel = 1; - (void)IsPotentiallyTruncated(V, BitWidth); + (void)for_each(E.Scalars, std::bind(IsPotentiallyTruncated, _1, + std::ref(BitWidth))); } else { // Several vectorized uses? Check if we can truncate it, otherwise - // exit. - if (ITE.UserTreeIndices.size() > 1 && - !IsPotentiallyTruncated(I, BitWidth)) + if (E.UserTreeIndices.size() > 1 && + !all_of(E.Scalars, std::bind(IsPotentiallyTruncated, _1, + std::ref(BitWidth)))) return false; bool NeedToExit = false; if (Checker && !AttemptCheckBitwidth(Checker, NeedToExit)) @@ -14266,26 +14264,22 @@ bool BoUpSLP::collectValuesToDemote( } ++MaxDepthLevel; - // Gather demoted constant operands. - for (unsigned Idx : seq(Start, End)) - if (isa(I->getOperand(Idx))) - DemotedConsts.try_emplace(I).first->getSecond().push_back(Idx); - // Record the value that we can demote. - ToDemote.push_back(V); + // Record the entry that we can demote. + ToDemote.push_back(E.Idx); return IsProfitableToDemote; }; - switch (I->getOpcode()) { + switch (E.getOpcode()) { // We can always demote truncations and extensions. Since truncations can // seed additional demotion, we save the truncated value. case Instruction::Trunc: if (IsProfitableToDemoteRoot) IsProfitableToDemote = true; - return TryProcessInstruction(I, *ITE, BitWidth); + return TryProcessInstruction(BitWidth); case Instruction::ZExt: case Instruction::SExt: IsProfitableToDemote = true; - return TryProcessInstruction(I, *ITE, BitWidth); + return TryProcessInstruction(BitWidth); // We can demote certain binary operations if we can demote both of their // operands. @@ -14295,112 +14289,128 @@ bool BoUpSLP::collectValuesToDemote( case Instruction::And: case Instruction::Or: case Instruction::Xor: { - return TryProcessInstruction(I, *ITE, BitWidth, - {I->getOperand(0), I->getOperand(1)}); + return TryProcessInstruction( + BitWidth, {getOperandEntry(&E, 0), getOperandEntry(&E, 1)}); } case Instruction::Shl: { // If we are truncating the result of this SHL, and if it's a shift of an // inrange amount, we can always perform a SHL in a smaller type. auto ShlChecker = [&](unsigned BitWidth, unsigned) { - KnownBits AmtKnownBits = computeKnownBits(I->getOperand(1), *DL); - return AmtKnownBits.getMaxValue().ult(BitWidth); + return all_of(E.Scalars, [&](Value *V) { + auto *I = cast(V); + KnownBits AmtKnownBits = computeKnownBits(I->getOperand(1), *DL); + return AmtKnownBits.getMaxValue().ult(BitWidth); + }); }; return TryProcessInstruction( - I, *ITE, BitWidth, {I->getOperand(0), I->getOperand(1)}, ShlChecker); + BitWidth, {getOperandEntry(&E, 0), getOperandEntry(&E, 1)}, ShlChecker); } case Instruction::LShr: { // If this is a truncate of a logical shr, we can truncate it to a smaller // lshr iff we know that the bits we would otherwise be shifting in are // already zeros. auto LShrChecker = [&](unsigned BitWidth, unsigned OrigBitWidth) { - KnownBits AmtKnownBits = computeKnownBits(I->getOperand(1), *DL); - APInt ShiftedBits = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); - return AmtKnownBits.getMaxValue().ult(BitWidth) && - MaskedValueIsZero(I->getOperand(0), ShiftedBits, - SimplifyQuery(*DL)); + return all_of(E.Scalars, [&](Value *V) { + auto *I = cast(V); + KnownBits AmtKnownBits = computeKnownBits(I->getOperand(1), *DL); + APInt ShiftedBits = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); + return AmtKnownBits.getMaxValue().ult(BitWidth) && + MaskedValueIsZero(I->getOperand(0), ShiftedBits, + SimplifyQuery(*DL)); + }); }; return TryProcessInstruction( - I, *ITE, BitWidth, {I->getOperand(0), I->getOperand(1)}, LShrChecker); + BitWidth, {getOperandEntry(&E, 0), getOperandEntry(&E, 1)}, + LShrChecker); } case Instruction::AShr: { // If this is a truncate of an arithmetic shr, we can truncate it to a // smaller ashr iff we know that all the bits from the sign bit of the // original type and the sign bit of the truncate type are similar. auto AShrChecker = [&](unsigned BitWidth, unsigned OrigBitWidth) { - KnownBits AmtKnownBits = computeKnownBits(I->getOperand(1), *DL); - unsigned ShiftedBits = OrigBitWidth - BitWidth; - return AmtKnownBits.getMaxValue().ult(BitWidth) && - ShiftedBits < - ComputeNumSignBits(I->getOperand(0), *DL, 0, AC, nullptr, DT); + return all_of(E.Scalars, [&](Value *V) { + auto *I = cast(V); + KnownBits AmtKnownBits = computeKnownBits(I->getOperand(1), *DL); + unsigned ShiftedBits = OrigBitWidth - BitWidth; + return AmtKnownBits.getMaxValue().ult(BitWidth) && + ShiftedBits < ComputeNumSignBits(I->getOperand(0), *DL, 0, AC, + nullptr, DT); + }); }; return TryProcessInstruction( - I, *ITE, BitWidth, {I->getOperand(0), I->getOperand(1)}, AShrChecker); + BitWidth, {getOperandEntry(&E, 0), getOperandEntry(&E, 1)}, + AShrChecker); } case Instruction::UDiv: case Instruction::URem: { // UDiv and URem can be truncated if all the truncated bits are zero. auto Checker = [&](unsigned BitWidth, unsigned OrigBitWidth) { assert(BitWidth <= OrigBitWidth && "Unexpected bitwidths!"); - APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); - return MaskedValueIsZero(I->getOperand(0), Mask, SimplifyQuery(*DL)) && - MaskedValueIsZero(I->getOperand(1), Mask, SimplifyQuery(*DL)); + return all_of(E.Scalars, [&](Value *V) { + auto *I = cast(V); + APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); + return MaskedValueIsZero(I->getOperand(0), Mask, SimplifyQuery(*DL)) && + MaskedValueIsZero(I->getOperand(1), Mask, SimplifyQuery(*DL)); + }); }; - return TryProcessInstruction(I, *ITE, BitWidth, - {I->getOperand(0), I->getOperand(1)}, Checker); + return TryProcessInstruction( + BitWidth, {getOperandEntry(&E, 0), getOperandEntry(&E, 1)}, Checker); } // We can demote selects if we can demote their true and false values. case Instruction::Select: { - Start = 1; - auto *SI = cast(I); - return TryProcessInstruction(I, *ITE, BitWidth, - {SI->getTrueValue(), SI->getFalseValue()}); + return TryProcessInstruction( + BitWidth, {getOperandEntry(&E, 1), getOperandEntry(&E, 2)}); } // We can demote phis if we can demote all their incoming operands. Note that // we don't need to worry about cycles since we ensure single use above. case Instruction::PHI: { - PHINode *PN = cast(I); - SmallVector Ops(PN->incoming_values().begin(), - PN->incoming_values().end()); - return TryProcessInstruction(I, *ITE, BitWidth, Ops); + const unsigned NumOps = E.getNumOperands(); + SmallVector Ops(NumOps); + transform(seq(0, NumOps), Ops.begin(), + std::bind(&BoUpSLP::getOperandEntry, this, &E, _1)); + + return TryProcessInstruction(BitWidth, Ops); } case Instruction::Call: { - auto *IC = dyn_cast(I); + auto *IC = dyn_cast(E.getMainOp()); if (!IC) break; Intrinsic::ID ID = getVectorIntrinsicIDForCall(IC, TLI); if (ID != Intrinsic::abs && ID != Intrinsic::smin && ID != Intrinsic::smax && ID != Intrinsic::umin && ID != Intrinsic::umax) break; - SmallVector Operands(1, I->getOperand(0)); + SmallVector Operands(1, getOperandEntry(&E, 0)); function_ref CallChecker; auto CompChecker = [&](unsigned BitWidth, unsigned OrigBitWidth) { assert(BitWidth <= OrigBitWidth && "Unexpected bitwidths!"); - if (ID == Intrinsic::umin || ID == Intrinsic::umax) { - APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); - return MaskedValueIsZero(I->getOperand(0), Mask, SimplifyQuery(*DL)) && - MaskedValueIsZero(I->getOperand(1), Mask, SimplifyQuery(*DL)); - } - assert((ID == Intrinsic::smin || ID == Intrinsic::smax) && - "Expected min/max intrinsics only."); - unsigned SignBits = OrigBitWidth - BitWidth; - return SignBits <= ComputeNumSignBits(I->getOperand(0), *DL, 0, AC, - nullptr, DT) && - SignBits <= - ComputeNumSignBits(I->getOperand(1), *DL, 0, AC, nullptr, DT); + return all_of(E.Scalars, [&](Value *V) { + auto *I = cast(V); + if (ID == Intrinsic::umin || ID == Intrinsic::umax) { + APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, BitWidth); + return MaskedValueIsZero(I->getOperand(0), Mask, + SimplifyQuery(*DL)) && + MaskedValueIsZero(I->getOperand(1), Mask, SimplifyQuery(*DL)); + } + assert((ID == Intrinsic::smin || ID == Intrinsic::smax) && + "Expected min/max intrinsics only."); + unsigned SignBits = OrigBitWidth - BitWidth; + return SignBits <= ComputeNumSignBits(I->getOperand(0), *DL, 0, AC, + nullptr, DT) && + SignBits <= ComputeNumSignBits(I->getOperand(1), *DL, 0, AC, + nullptr, DT); + }); }; - End = 1; if (ID != Intrinsic::abs) { - Operands.push_back(I->getOperand(1)); - End = 2; + Operands.push_back(getOperandEntry(&E, 1)); CallChecker = CompChecker; } InstructionCost BestCost = std::numeric_limits::max(); unsigned BestBitWidth = BitWidth; - unsigned VF = ITE->Scalars.size(); + unsigned VF = E.Scalars.size(); // Choose the best bitwidth based on cost estimations. auto Checker = [&](unsigned BitWidth, unsigned) { unsigned MinBW = PowerOf2Ceil(BitWidth); @@ -14419,7 +14429,7 @@ bool BoUpSLP::collectValuesToDemote( [[maybe_unused]] bool NeedToExit; (void)AttemptCheckBitwidth(Checker, NeedToExit); BitWidth = BestBitWidth; - return TryProcessInstruction(I, *ITE, BitWidth, Operands, CallChecker); + return TryProcessInstruction(BitWidth, Operands, CallChecker); } // Otherwise, conservatively give up. @@ -14473,26 +14483,27 @@ void BoUpSLP::computeMinimumValueSizes() { ++NodeIdx; } - // Analyzed in reduction already and not profitable - exit. + // Analyzed the reduction already and not profitable - exit. if (AnalyzedMinBWVals.contains(VectorizableTree[NodeIdx]->Scalars.front())) return; - SmallVector ToDemote; - DenseMap> DemotedConsts; - auto ComputeMaxBitWidth = [&](ArrayRef TreeRoot, unsigned VF, - bool IsTopRoot, bool IsProfitableToDemoteRoot, - unsigned Opcode, unsigned Limit, - bool IsTruncRoot, bool IsSignedCmp) { + SmallVector ToDemote; + auto ComputeMaxBitWidth = [&](const TreeEntry &E, bool IsTopRoot, + bool IsProfitableToDemoteRoot, unsigned Opcode, + unsigned Limit, bool IsTruncRoot, + bool IsSignedCmp) { ToDemote.clear(); - auto *TreeRootIT = dyn_cast(TreeRoot[0]->getType()); + unsigned VF = E.getVectorFactor(); + auto *TreeRootIT = dyn_cast(E.Scalars.front()->getType()); if (!TreeRootIT || !Opcode) return 0u; - if (AnalyzedMinBWVals.contains(TreeRoot.front())) + if (any_of(E.Scalars, + [&](Value *V) { return AnalyzedMinBWVals.contains(V); })) return 0u; - unsigned NumParts = TTI->getNumberOfParts( - FixedVectorType::get(TreeRoot.front()->getType(), VF)); + unsigned NumParts = + TTI->getNumberOfParts(FixedVectorType::get(TreeRootIT, VF)); // The maximum bit width required to represent all the values that can be // demoted without loss of precision. It would be safe to truncate the roots @@ -14505,14 +14516,14 @@ void BoUpSLP::computeMinimumValueSizes() { // True. // Determine if the sign bit of all the roots is known to be zero. If not, // IsKnownPositive is set to False. - bool IsKnownPositive = !IsSignedCmp && all_of(TreeRoot, [&](Value *R) { + bool IsKnownPositive = !IsSignedCmp && all_of(E.Scalars, [&](Value *R) { KnownBits Known = computeKnownBits(R, *DL); return Known.isNonNegative(); }); // We first check if all the bits of the roots are demanded. If they're not, // we can truncate the roots to this narrower type. - for (auto *Root : TreeRoot) { + for (Value *Root : E.Scalars) { unsigned NumSignBits = ComputeNumSignBits(Root, *DL, 0, AC, nullptr, DT); TypeSize NumTypeBits = DL->getTypeSizeInBits(Root->getType()); unsigned BitWidth1 = NumTypeBits - NumSignBits; @@ -14557,23 +14568,22 @@ void BoUpSLP::computeMinimumValueSizes() { // Conservatively determine if we can actually truncate the roots of the // expression. Collect the values that can be demoted in ToDemote and // additional roots that require investigating in Roots. - for (auto *Root : TreeRoot) { - DenseSet Visited; - unsigned MaxDepthLevel = IsTruncRoot ? Limit : 1; - bool NeedToDemote = IsProfitableToDemote; - - if (!collectValuesToDemote(Root, IsProfitableToDemoteRoot, MaxBitWidth, - ToDemote, DemotedConsts, Visited, - MaxDepthLevel, NeedToDemote, IsTruncRoot) || - (MaxDepthLevel <= Limit && - !(((Opcode == Instruction::SExt || Opcode == Instruction::ZExt) && - (!IsTopRoot || !(IsStoreOrInsertElt || UserIgnoreList) || - DL->getTypeSizeInBits(Root->getType()) / - DL->getTypeSizeInBits( - cast(Root)->getOperand(0)->getType()) > - 2))))) - return 0u; - } + DenseSet Visited; + unsigned MaxDepthLevel = IsTruncRoot ? Limit : 1; + bool NeedToDemote = IsProfitableToDemote; + + if (!collectValuesToDemote(E, IsProfitableToDemoteRoot, MaxBitWidth, + ToDemote, Visited, MaxDepthLevel, NeedToDemote, + IsTruncRoot) || + (MaxDepthLevel <= Limit && + !(((Opcode == Instruction::SExt || Opcode == Instruction::ZExt) && + (!IsTopRoot || !(IsStoreOrInsertElt || UserIgnoreList) || + DL->getTypeSizeInBits(TreeRootIT) / + DL->getTypeSizeInBits(cast(E.Scalars.front()) + ->getOperand(0) + ->getType()) > + 2))))) + return 0u; // Round MaxBitWidth up to the next power-of-two. MaxBitWidth = bit_ceil(MaxBitWidth); @@ -14624,8 +14634,8 @@ void BoUpSLP::computeMinimumValueSizes() { VectorizableTree.front()->Scalars.front()->getType())) Limit = 3; unsigned MaxBitWidth = ComputeMaxBitWidth( - TreeRoot, VectorizableTree[NodeIdx]->getVectorFactor(), IsTopRoot, - IsProfitableToDemoteRoot, Opcode, Limit, IsTruncRoot, IsSignedCmp); + *VectorizableTree[NodeIdx].get(), IsTopRoot, IsProfitableToDemoteRoot, + Opcode, Limit, IsTruncRoot, IsSignedCmp); if (ReductionBitWidth != 0 && (IsTopRoot || !RootDemotes.empty())) { if (MaxBitWidth != 0 && ReductionBitWidth < MaxBitWidth) ReductionBitWidth = bit_ceil(MaxBitWidth); @@ -14634,13 +14644,15 @@ void BoUpSLP::computeMinimumValueSizes() { } for (unsigned Idx : RootDemotes) { - Value *V = VectorizableTree[Idx]->Scalars.front(); - uint32_t OrigBitWidth = DL->getTypeSizeInBits(V->getType()); - if (OrigBitWidth > MaxBitWidth) { - APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, MaxBitWidth); - if (MaskedValueIsZero(V, Mask, SimplifyQuery(*DL))) - ToDemote.push_back(V); - } + if (all_of(VectorizableTree[Idx]->Scalars, [&](Value *V) { + uint32_t OrigBitWidth = DL->getTypeSizeInBits(V->getType()); + if (OrigBitWidth > MaxBitWidth) { + APInt Mask = APInt::getBitsSetFrom(OrigBitWidth, MaxBitWidth); + return MaskedValueIsZero(V, Mask, SimplifyQuery(*DL)); + } + return false; + })) + ToDemote.push_back(Idx); } RootDemotes.clear(); IsTopRoot = false; @@ -14687,9 +14699,8 @@ void BoUpSLP::computeMinimumValueSizes() { // Finally, map the values we can demote to the maximum bit with we // computed. - for (Value *Scalar : ToDemote) { - TreeEntry *TE = getTreeEntry(Scalar); - assert(TE && "Expected vectorized scalar."); + for (unsigned Idx : ToDemote) { + TreeEntry *TE = VectorizableTree[Idx].get(); if (MinBWs.contains(TE)) continue; bool IsSigned = TE->getOpcode() == Instruction::SExt || @@ -14697,22 +14708,6 @@ void BoUpSLP::computeMinimumValueSizes() { return !isKnownNonNegative(R, SimplifyQuery(*DL)); }); MinBWs.try_emplace(TE, MaxBitWidth, IsSigned); - const auto *I = cast(Scalar); - auto DCIt = DemotedConsts.find(I); - if (DCIt != DemotedConsts.end()) { - for (unsigned Idx : DCIt->getSecond()) { - // Check that all instructions operands are demoted. - const TreeEntry *CTE = getOperandEntry(TE, Idx); - if (all_of(TE->Scalars, - [&](Value *V) { - auto SIt = DemotedConsts.find(cast(V)); - return SIt != DemotedConsts.end() && - is_contained(SIt->getSecond(), Idx); - }) || - all_of(CTE->Scalars, IsaPred)) - MinBWs.try_emplace(CTE, MaxBitWidth, IsSigned); - } - } } } } diff --git a/llvm/test/Analysis/ScalarEvolution/pr87798.ll b/llvm/test/Analysis/ScalarEvolution/pr87798.ll new file mode 100644 index 0000000000000..acd445993e47b --- /dev/null +++ b/llvm/test/Analysis/ScalarEvolution/pr87798.ll @@ -0,0 +1,68 @@ +; NOTE: Assertions have been autogenerated by utils/update_analyze_test_checks.py UTC_ARGS: --version 4 +; RUN: opt -disable-output -passes='print' -verify-scev < %s 2>&1 | FileCheck %s + +target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-i128:128-f80:128-n8:16:32:64-S128-ni:1-p2:32:8:8:32-ni:2" +target triple = "x86_64-unknown-linux-gnu" + +; print is used to compute SCEVs for all values in the +; function. +; We should not crash on multiplicative inverse called within SCEV's binomial +; coefficient function. + +define i32 @pr87798() { +; CHECK-LABEL: 'pr87798' +; CHECK-NEXT: Classifying expressions for: @pr87798 +; CHECK-NEXT: %phi = phi i32 [ 0, %bb ], [ %add4, %bb1 ] +; CHECK-NEXT: --> {0,+,0,+,0,+,2,+,3}<%bb1> U: full-set S: full-set Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %phi2 = phi i32 [ 0, %bb ], [ %add, %bb1 ] +; CHECK-NEXT: --> {0,+,0,+,1}<%bb1> U: full-set S: full-set Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %phi3 = phi i32 [ 0, %bb ], [ %add5, %bb1 ] +; CHECK-NEXT: --> {0,+,1}<%bb1> U: [0,1) S: [0,1) Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %add = add i32 %phi2, %phi3 +; CHECK-NEXT: --> {0,+,1,+,1}<%bb1> U: full-set S: full-set Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %mul = mul i32 %phi2, %phi3 +; CHECK-NEXT: --> {0,+,0,+,2,+,3}<%bb1> U: full-set S: full-set Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %add4 = add i32 %mul, %phi +; CHECK-NEXT: --> {0,+,0,+,2,+,5,+,3}<%bb1> U: full-set S: full-set Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %and = and i32 %phi, 1 +; CHECK-NEXT: --> (zext i1 {false,+,false,+,false,+,false,+,true}<%bb1> to i32) U: [0,2) S: [0,2) Exits: 0 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %add5 = add i32 %phi3, 1 +; CHECK-NEXT: --> {1,+,1}<%bb1> U: [1,2) S: [1,2) Exits: 1 LoopDispositions: { %bb1: Computable } +; CHECK-NEXT: %phi9 = phi i32 [ %and, %bb1 ] +; CHECK-NEXT: --> (zext i1 {false,+,false,+,false,+,false,+,true}<%bb1> to i32) U: [0,2) S: [0,2) --> 0 U: [0,1) S: [0,1) +; CHECK-NEXT: %zext = zext i32 %phi9 to i64 +; CHECK-NEXT: --> poison U: full-set S: full-set +; CHECK-NEXT: Determining loop execution counts for: @pr87798 +; CHECK-NEXT: Loop %loop: Unpredictable backedge-taken count. +; CHECK-NEXT: Loop %loop: Unpredictable constant max backedge-taken count. +; CHECK-NEXT: Loop %loop: Unpredictable symbolic max backedge-taken count. +; CHECK-NEXT: Loop %bb1: backedge-taken count is i1 false +; CHECK-NEXT: Loop %bb1: constant max backedge-taken count is i1 false +; CHECK-NEXT: Loop %bb1: symbolic max backedge-taken count is i1 false +; CHECK-NEXT: Loop %bb1: Trip multiple is 1 +; +bb: + br label %bb1 + +bb1: ; preds = %bb1, %bb + %phi = phi i32 [ 0, %bb ], [ %add4, %bb1 ] + %phi2 = phi i32 [ 0, %bb ], [ %add, %bb1 ] + %phi3 = phi i32 [ 0, %bb ], [ %add5, %bb1 ] + %add = add i32 %phi2, %phi3 + %mul = mul i32 %phi2, %phi3 + %add4 = add i32 %mul, %phi + %and = and i32 %phi, 1 + %add5 = add i32 %phi3, 1 + br i1 true, label %preheader, label %bb1 + +preheader: ; preds = %bb1 + %phi9 = phi i32 [ %and, %bb1 ] + br label %loop + +loop: ; preds = %preheader, %loop + br label %loop + +bb7: ; No predecessors! + %zext = zext i32 %phi9 to i64 + ret i32 0 +} diff --git a/llvm/test/Transforms/PhaseOrdering/AArch64/hoist-runtime-checks.ll b/llvm/test/Transforms/PhaseOrdering/AArch64/hoist-runtime-checks.ll index c6c9a52167d54..a140e17a0dd15 100644 --- a/llvm/test/Transforms/PhaseOrdering/AArch64/hoist-runtime-checks.ll +++ b/llvm/test/Transforms/PhaseOrdering/AArch64/hoist-runtime-checks.ll @@ -91,8 +91,151 @@ for.end: ; preds = %for.cond.cleanup ret i32 %9 } +%"class.std::__1::span" = type { ptr, i64 } +%"class.std::__1::__wrap_iter" = type { ptr } + +define dso_local noundef i32 @sum_prefix_with_sum(ptr %s.coerce0, i64 %s.coerce1, i64 noundef %n) { +; CHECK-LABEL: define dso_local noundef i32 @sum_prefix_with_sum( +; CHECK-SAME: ptr nocapture readonly [[S_COERCE0:%.*]], i64 [[S_COERCE1:%.*]], i64 noundef [[N:%.*]]) local_unnamed_addr #[[ATTR0]] { +; CHECK-NEXT: entry: +; CHECK-NEXT: [[CMP5_NOT:%.*]] = icmp eq i64 [[N]], 0 +; CHECK-NEXT: br i1 [[CMP5_NOT]], label [[FOR_COND_CLEANUP:%.*]], label [[FOR_BODY_PREHEADER:%.*]] +; CHECK: for.body.preheader: +; CHECK-NEXT: [[TMP0:%.*]] = add i64 [[N]], -1 +; CHECK-NEXT: [[DOTNOT_NOT:%.*]] = icmp ult i64 [[TMP0]], [[S_COERCE1]] +; CHECK-NEXT: br label [[FOR_BODY:%.*]] +; CHECK: for.cond.cleanup: +; CHECK-NEXT: [[RET_0_LCSSA:%.*]] = phi i32 [ 0, [[ENTRY:%.*]] ], [ [[ADD:%.*]], [[SPAN_CHECKED_ACCESS_EXIT:%.*]] ] +; CHECK-NEXT: ret i32 [[RET_0_LCSSA]] +; CHECK: for.body: +; CHECK-NEXT: [[I_07:%.*]] = phi i64 [ [[INC:%.*]], [[SPAN_CHECKED_ACCESS_EXIT]] ], [ 0, [[FOR_BODY_PREHEADER]] ] +; CHECK-NEXT: [[RET_06:%.*]] = phi i32 [ [[ADD]], [[SPAN_CHECKED_ACCESS_EXIT]] ], [ 0, [[FOR_BODY_PREHEADER]] ] +; CHECK-NEXT: br i1 [[DOTNOT_NOT]], label [[SPAN_CHECKED_ACCESS_EXIT]], label [[COND_FALSE_I:%.*]], !prof [[PROF0:![0-9]+]] +; CHECK: cond.false.i: +; CHECK-NEXT: tail call void @llvm.trap() +; CHECK-NEXT: unreachable +; CHECK: span_checked_access.exit: +; CHECK-NEXT: [[ARRAYIDX_I:%.*]] = getelementptr inbounds i32, ptr [[S_COERCE0]], i64 [[I_07]] +; CHECK-NEXT: [[TMP7:%.*]] = load i32, ptr [[ARRAYIDX_I]], align 4 +; CHECK-NEXT: [[ADD]] = add nsw i32 [[TMP7]], [[RET_06]] +; CHECK-NEXT: [[INC]] = add nuw i64 [[I_07]], 1 +; CHECK-NEXT: [[EXITCOND_NOT:%.*]] = icmp eq i64 [[INC]], [[N]] +; CHECK-NEXT: br i1 [[EXITCOND_NOT]], label [[FOR_COND_CLEANUP]], label [[FOR_BODY]] +; +entry: + %s = alloca %"class.std::__1::span", align 8 + %n.addr = alloca i64, align 8 + %ret = alloca i32, align 4 + %i = alloca i64, align 8 + %0 = getelementptr inbounds { ptr, i64 }, ptr %s, i32 0, i32 0 + store ptr %s.coerce0, ptr %0, align 8 + %1 = getelementptr inbounds { ptr, i64 }, ptr %s, i32 0, i32 1 + store i64 %s.coerce1, ptr %1, align 8 + store i64 %n, ptr %n.addr, align 8 + call void @llvm.lifetime.start.p0(i64 4, ptr %ret) #7 + store i32 0, ptr %ret, align 4 + call void @llvm.lifetime.start.p0(i64 8, ptr %i) #7 + store i64 0, ptr %i, align 8 + br label %for.cond + +for.cond: ; preds = %for.inc, %entry + %2 = load i64, ptr %i, align 8 + %3 = load i64, ptr %n.addr, align 8 + %cmp = icmp ult i64 %2, %3 + br i1 %cmp, label %for.body, label %for.cond.cleanup + +for.cond.cleanup: ; preds = %for.cond + call void @llvm.lifetime.end.p0(i64 8, ptr %i) #7 + br label %for.end + +for.body: ; preds = %for.cond + %4 = load i64, ptr %i, align 8 + %call = call noundef nonnull align 4 dereferenceable(4) ptr @span_checked_access(ptr noundef nonnull align 8 dereferenceable(16) %s, i64 noundef %4) #7 + %5 = load i32, ptr %call, align 4 + %6 = load i32, ptr %ret, align 4 + %add = add nsw i32 %6, %5 + store i32 %add, ptr %ret, align 4 + br label %for.inc + +for.inc: ; preds = %for.body + %7 = load i64, ptr %i, align 8 + %inc = add i64 %7, 1 + store i64 %inc, ptr %i, align 8 + br label %for.cond + +for.end: ; preds = %for.cond.cleanup + %8 = load i32, ptr %ret, align 4 + call void @llvm.lifetime.end.p0(i64 4, ptr %ret) + ret i32 %8 +} + +define hidden noundef nonnull align 4 dereferenceable(4) ptr @span_checked_access(ptr noundef nonnull align 8 dereferenceable(16) %this, i64 noundef %__idx) { +; CHECK-LABEL: define hidden noundef nonnull align 4 dereferenceable(4) ptr @span_checked_access( +; CHECK-SAME: ptr nocapture noundef nonnull readonly align 8 dereferenceable(16) [[THIS:%.*]], i64 noundef [[__IDX:%.*]]) local_unnamed_addr #[[ATTR0]] { +; CHECK-NEXT: entry: +; CHECK-NEXT: [[__SIZE__I:%.*]] = getelementptr inbounds i8, ptr [[THIS]], i64 8 +; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[__SIZE__I]], align 8 +; CHECK-NEXT: [[CMP:%.*]] = icmp ugt i64 [[TMP0]], [[__IDX]] +; CHECK-NEXT: br i1 [[CMP]], label [[COND_END:%.*]], label [[COND_FALSE:%.*]], !prof [[PROF0]] +; CHECK: cond.false: +; CHECK-NEXT: tail call void @llvm.trap() +; CHECK-NEXT: unreachable +; CHECK: cond.end: +; CHECK-NEXT: [[TMP1:%.*]] = load ptr, ptr [[THIS]], align 8 +; CHECK-NEXT: [[ARRAYIDX:%.*]] = getelementptr inbounds i32, ptr [[TMP1]], i64 [[__IDX]] +; CHECK-NEXT: ret ptr [[ARRAYIDX]] +; +entry: + %this.addr = alloca ptr, align 8 + %__idx.addr = alloca i64, align 8 + store ptr %this, ptr %this.addr, align 8 + store i64 %__idx, ptr %__idx.addr, align 8 + %this1 = load ptr, ptr %this.addr, align 8 + %0 = load i64, ptr %__idx.addr, align 8 + %call = call noundef i64 @span_access(ptr noundef nonnull align 8 dereferenceable(16) %this1) + %cmp = icmp ult i64 %0, %call + %conv = zext i1 %cmp to i64 + %expval = call i64 @llvm.expect.i64(i64 %conv, i64 1) + %tobool = icmp ne i64 %expval, 0 + br i1 %tobool, label %cond.true, label %cond.false + +cond.true: ; preds = %entry + br label %cond.end + +cond.false: ; preds = %entry + call void @llvm.trap() + br label %cond.end + +cond.end: ; preds = %cond.false, %cond.true + %__data_ = getelementptr inbounds %"class.std::__1::span", ptr %this1, i32 0, i32 0 + %1 = load ptr, ptr %__data_, align 8 + %2 = load i64, ptr %__idx.addr, align 8 + %arrayidx = getelementptr inbounds i32, ptr %1, i64 %2 + ret ptr %arrayidx +} + +define hidden noundef i64 @span_access(ptr noundef nonnull align 8 dereferenceable(16) %this) { +; CHECK-LABEL: define hidden noundef i64 @span_access( +; CHECK-SAME: ptr nocapture noundef nonnull readonly align 8 dereferenceable(16) [[THIS:%.*]]) local_unnamed_addr #[[ATTR1:[0-9]+]] { +; CHECK-NEXT: entry: +; CHECK-NEXT: [[__SIZE_:%.*]] = getelementptr inbounds i8, ptr [[THIS]], i64 8 +; CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[__SIZE_]], align 8 +; CHECK-NEXT: ret i64 [[TMP0]] +; +entry: + %this.addr = alloca ptr, align 8 + store ptr %this, ptr %this.addr, align 8 + %this1 = load ptr, ptr %this.addr, align 8 + %__size_ = getelementptr inbounds %"class.std::__1::span", ptr %this1, i32 0, i32 1 + %0 = load i64, ptr %__size_, align 8 + ret i64 %0 +} + declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture) declare void @llvm.trap() declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture) +;. +; CHECK: [[PROF0]] = !{!"branch_weights", i32 2000, i32 1} +;. diff --git a/mlir/include/mlir/IR/OpDefinition.h b/mlir/include/mlir/IR/OpDefinition.h index c177ae3594d11..2d1dee2303e8f 100644 --- a/mlir/include/mlir/IR/OpDefinition.h +++ b/mlir/include/mlir/IR/OpDefinition.h @@ -1965,7 +1965,7 @@ class Op : public OpState, public Traits... { if constexpr (!hasProperties()) return getEmptyProperties(); return *getOperation() - ->getPropertiesStorage() + ->getPropertiesStorageUnsafe() .template as *>(); } diff --git a/mlir/include/mlir/IR/Operation.h b/mlir/include/mlir/IR/Operation.h index 3ffd3517fe5a6..c52a6fcac10c1 100644 --- a/mlir/include/mlir/IR/Operation.h +++ b/mlir/include/mlir/IR/Operation.h @@ -895,8 +895,7 @@ class alignas(8) Operation final /// Returns the properties storage. OpaqueProperties getPropertiesStorage() { if (propertiesStorageSize) - return { - reinterpret_cast(getTrailingObjects())}; + return getPropertiesStorageUnsafe(); return {nullptr}; } OpaqueProperties getPropertiesStorage() const { @@ -905,6 +904,12 @@ class alignas(8) Operation final getTrailingObjects()))}; return {nullptr}; } + /// Returns the properties storage without checking whether properties are + /// present. + OpaqueProperties getPropertiesStorageUnsafe() { + return { + reinterpret_cast(getTrailingObjects())}; + } /// Return the properties converted to an attribute. /// This is expensive, and mostly useful when dealing with unregistered diff --git a/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp b/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp index 9c3c4d96a301e..0aa1de5fa5d9a 100644 --- a/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp +++ b/mlir/lib/Conversion/ComplexToStandard/ComplexToStandard.cpp @@ -570,37 +570,39 @@ struct Log1pOpConversion : public OpConversionPattern { ConversionPatternRewriter &rewriter) const override { auto type = cast(adaptor.getComplex().getType()); auto elementType = cast(type.getElementType()); - arith::FastMathFlagsAttr fmf = op.getFastMathFlagsAttr(); + arith::FastMathFlags fmf = op.getFastMathFlagsAttr().getValue(); mlir::ImplicitLocOpBuilder b(op.getLoc(), rewriter); - Value real = b.create(elementType, adaptor.getComplex()); - Value imag = b.create(elementType, adaptor.getComplex()); + Value real = b.create(adaptor.getComplex()); + Value imag = b.create(adaptor.getComplex()); Value half = b.create(elementType, b.getFloatAttr(elementType, 0.5)); Value one = b.create(elementType, b.getFloatAttr(elementType, 1)); - Value two = b.create(elementType, - b.getFloatAttr(elementType, 2)); - - // log1p(a+bi) = .5*log((a+1)^2+b^2) + i*atan2(b, a + 1) - // log((a+1)+bi) = .5*log(a*a + 2*a + 1 + b*b) + i*atan2(b, a+1) - // log((a+1)+bi) = .5*log1p(a*a + 2*a + b*b) + i*atan2(b, a+1) - Value sumSq = b.create(real, real, fmf.getValue()); - sumSq = b.create( - sumSq, b.create(real, two, fmf.getValue()), - fmf.getValue()); - sumSq = b.create( - sumSq, b.create(imag, imag, fmf.getValue()), - fmf.getValue()); - Value logSumSq = - b.create(elementType, sumSq, fmf.getValue()); - Value resultReal = b.create(logSumSq, half, fmf.getValue()); - - Value realPlusOne = b.create(real, one, fmf.getValue()); - - Value resultImag = - b.create(elementType, imag, realPlusOne, fmf.getValue()); + Value realPlusOne = b.create(real, one, fmf); + Value absRealPlusOne = b.create(realPlusOne, fmf); + Value absImag = b.create(imag, fmf); + + Value maxAbs = b.create(absRealPlusOne, absImag, fmf); + Value minAbs = b.create(absRealPlusOne, absImag, fmf); + + Value maxAbsOfRealPlusOneAndImagMinusOne = b.create( + b.create(arith::CmpFPredicate::OGT, realPlusOne, absImag, + fmf), + real, b.create(maxAbs, one, fmf)); + Value minMaxRatio = b.create(minAbs, maxAbs, fmf); + Value logOfMaxAbsOfRealPlusOneAndImag = + b.create(maxAbsOfRealPlusOneAndImagMinusOne, fmf); + Value logOfSqrtPart = b.create( + b.create(minMaxRatio, minMaxRatio, fmf), fmf); + Value r = b.create( + b.create(half, logOfSqrtPart, fmf), + logOfMaxAbsOfRealPlusOneAndImag, fmf); + Value resultReal = b.create( + b.create(arith::CmpFPredicate::UNO, r, r, fmf), minAbs, + r); + Value resultImag = b.create(imag, realPlusOne, fmf); rewriter.replaceOpWithNewOp(op, type, resultReal, resultImag); return success(); diff --git a/mlir/test/Conversion/ComplexToStandard/convert-to-standard.mlir b/mlir/test/Conversion/ComplexToStandard/convert-to-standard.mlir index f5d9499eadda4..43918904a09f4 100644 --- a/mlir/test/Conversion/ComplexToStandard/convert-to-standard.mlir +++ b/mlir/test/Conversion/ComplexToStandard/convert-to-standard.mlir @@ -300,15 +300,22 @@ func.func @complex_log1p(%arg: complex) -> complex { // CHECK: %[[IMAG:.*]] = complex.im %[[ARG]] : complex // CHECK: %[[ONE_HALF:.*]] = arith.constant 5.000000e-01 : f32 // CHECK: %[[ONE:.*]] = arith.constant 1.000000e+00 : f32 -// CHECK: %[[TWO:.*]] = arith.constant 2.000000e+00 : f32 -// CHECK: %[[SQ_SUM_0:.*]] = arith.mulf %[[REAL]], %[[REAL]] : f32 -// CHECK: %[[TWO_REAL:.*]] = arith.mulf %[[REAL]], %[[TWO]] : f32 -// CHECK: %[[SQ_SUM_1:.*]] = arith.addf %[[SQ_SUM_0]], %[[TWO_REAL]] : f32 -// CHECK: %[[SQ_IMAG:.*]] = arith.mulf %[[IMAG]], %[[IMAG]] : f32 -// CHECK: %[[SQ_SUM_2:.*]] = arith.addf %[[SQ_SUM_1]], %[[SQ_IMAG]] : f32 -// CHECK: %[[LOG_SQ_SUM:.*]] = math.log1p %[[SQ_SUM_2]] : f32 -// CHECK: %[[RESULT_REAL:.*]] = arith.mulf %[[LOG_SQ_SUM]], %[[ONE_HALF]] : f32 // CHECK: %[[REAL_PLUS_ONE:.*]] = arith.addf %[[REAL]], %[[ONE]] : f32 +// CHECK: %[[ABS_REAL_PLUS_ONE:.*]] = math.absf %[[REAL_PLUS_ONE]] : f32 +// CHECK: %[[ABS_IMAG:.*]] = math.absf %[[IMAG]] : f32 +// CHECK: %[[MAX:.*]] = arith.maximumf %[[ABS_REAL_PLUS_ONE]], %[[ABS_IMAG]] : f32 +// CHECK: %[[MIN:.*]] = arith.minimumf %[[ABS_REAL_PLUS_ONE]], %[[ABS_IMAG]] : f32 +// CHECK: %[[CMPF:.*]] = arith.cmpf ogt, %[[REAL_PLUS_ONE]], %[[ABS_IMAG]] : f32 +// CHECK: %[[MAX_MINUS_ONE:.*]] = arith.subf %[[MAX]], %cst_0 : f32 +// CHECK: %[[SELECT:.*]] = arith.select %[[CMPF]], %0, %[[MAX_MINUS_ONE]] : f32 +// CHECK: %[[MIN_MAX_RATIO:.*]] = arith.divf %[[MIN]], %[[MAX]] : f32 +// CHECK: %[[LOG_1:.*]] = math.log1p %[[SELECT]] : f32 +// CHECK: %[[RATIO_SQ:.*]] = arith.mulf %[[MIN_MAX_RATIO]], %[[MIN_MAX_RATIO]] : f32 +// CHECK: %[[LOG_SQ:.*]] = math.log1p %[[RATIO_SQ]] : f32 +// CHECK: %[[HALF_LOG_SQ:.*]] = arith.mulf %cst, %[[LOG_SQ]] : f32 +// CHECK: %[[R:.*]] = arith.addf %[[HALF_LOG_SQ]], %[[LOG_1]] : f32 +// CHECK: %[[ISNAN:.*]] = arith.cmpf uno, %[[R]], %[[R]] : f32 +// CHECK: %[[RESULT_REAL:.*]] = arith.select %[[ISNAN]], %[[MIN]], %[[R]] : f32 // CHECK: %[[RESULT_IMAG:.*]] = math.atan2 %[[IMAG]], %[[REAL_PLUS_ONE]] : f32 // CHECK: %[[RESULT:.*]] = complex.create %[[RESULT_REAL]], %[[RESULT_IMAG]] : complex // CHECK: return %[[RESULT]] : complex @@ -963,15 +970,22 @@ func.func @complex_log1p_with_fmf(%arg: complex) -> complex { // CHECK: %[[IMAG:.*]] = complex.im %[[ARG]] : complex // CHECK: %[[ONE_HALF:.*]] = arith.constant 5.000000e-01 : f32 // CHECK: %[[ONE:.*]] = arith.constant 1.000000e+00 : f32 -// CHECK: %[[TWO:.*]] = arith.constant 2.000000e+00 : f32 -// CHECK: %[[SQ_SUM_0:.*]] = arith.mulf %[[REAL]], %[[REAL]] fastmath : f32 -// CHECK: %[[TWO_REAL:.*]] = arith.mulf %[[REAL]], %[[TWO]] fastmath : f32 -// CHECK: %[[SQ_SUM_1:.*]] = arith.addf %[[SQ_SUM_0]], %[[TWO_REAL]] fastmath : f32 -// CHECK: %[[SQ_IMAG:.*]] = arith.mulf %[[IMAG]], %[[IMAG]] fastmath : f32 -// CHECK: %[[SQ_SUM_2:.*]] = arith.addf %[[SQ_SUM_1]], %[[SQ_IMAG]] fastmath : f32 -// CHECK: %[[LOG_SQ_SUM:.*]] = math.log1p %[[SQ_SUM_2]] fastmath : f32 -// CHECK: %[[RESULT_REAL:.*]] = arith.mulf %[[LOG_SQ_SUM]], %[[ONE_HALF]] fastmath : f32 -// CHECK: %[[REAL_PLUS_ONE:.*]] = arith.addf %[[REAL]], %[[ONE]] fastmath : f32 +// CHECK: %[[REAL_PLUS_ONE:.*]] = arith.addf %[[REAL]], %[[ONE]] fastmath : f32 +// CHECK: %[[ABS_REAL_PLUS_ONE:.*]] = math.absf %[[REAL_PLUS_ONE]] fastmath : f32 +// CHECK: %[[ABS_IMAG:.*]] = math.absf %[[IMAG]] fastmath : f32 +// CHECK: %[[MAX:.*]] = arith.maximumf %[[ABS_REAL_PLUS_ONE]], %[[ABS_IMAG]] fastmath : f32 +// CHECK: %[[MIN:.*]] = arith.minimumf %[[ABS_REAL_PLUS_ONE]], %[[ABS_IMAG]] fastmath : f32 +// CHECK: %[[CMPF:.*]] = arith.cmpf ogt, %[[REAL_PLUS_ONE]], %[[ABS_IMAG]] fastmath : f32 +// CHECK: %[[MAX_MINUS_ONE:.*]] = arith.subf %[[MAX]], %cst_0 fastmath : f32 +// CHECK: %[[SELECT:.*]] = arith.select %[[CMPF]], %0, %[[MAX_MINUS_ONE]] : f32 +// CHECK: %[[MIN_MAX_RATIO:.*]] = arith.divf %[[MIN]], %[[MAX]] fastmath : f32 +// CHECK: %[[LOG_1:.*]] = math.log1p %[[SELECT]] fastmath : f32 +// CHECK: %[[RATIO_SQ:.*]] = arith.mulf %[[MIN_MAX_RATIO]], %[[MIN_MAX_RATIO]] fastmath : f32 +// CHECK: %[[LOG_SQ:.*]] = math.log1p %[[RATIO_SQ]] fastmath : f32 +// CHECK: %[[HALF_LOG_SQ:.*]] = arith.mulf %cst, %[[LOG_SQ]] fastmath : f32 +// CHECK: %[[R:.*]] = arith.addf %[[HALF_LOG_SQ]], %[[LOG_1]] fastmath : f32 +// CHECK: %[[ISNAN:.*]] = arith.cmpf uno, %[[R]], %[[R]] fastmath : f32 +// CHECK: %[[RESULT_REAL:.*]] = arith.select %[[ISNAN]], %[[MIN]], %[[R]] : f32 // CHECK: %[[RESULT_IMAG:.*]] = math.atan2 %[[IMAG]], %[[REAL_PLUS_ONE]] fastmath : f32 // CHECK: %[[RESULT:.*]] = complex.create %[[RESULT_REAL]], %[[RESULT_IMAG]] : complex // CHECK: return %[[RESULT]] : complex