Date: Sat, 2 Jan 2021 21:20:01 GMT From: Dimitry Andric <dim@FreeBSD.org> To: src-committers@FreeBSD.org, dev-commits-src-all@FreeBSD.org Subject: git: 1de062e1e530 - vendor/llvm-project/release-11.x - Vendor import of llvm-project branch release/11.x llvmorg-11.0.1-rc2-0-g43ff75f2c3f (aka 11.0.1 rc2). Message-ID: <202101022120.102LK1tB066776@gitrepo.freebsd.org>
next in thread | raw e-mail | index | archive | help
The branch vendor/llvm-project/release-11.x has been updated by dim: URL: https://cgit.FreeBSD.org/src/commit/?id=1de062e1e530408f5a06466742b26722c230c024 commit 1de062e1e530408f5a06466742b26722c230c024 Author: Dimitry Andric <dim@FreeBSD.org> AuthorDate: 2021-01-02 21:13:33 +0000 Commit: Dimitry Andric <dim@FreeBSD.org> CommitDate: 2021-01-02 21:13:33 +0000 Vendor import of llvm-project branch release/11.x llvmorg-11.0.1-rc2-0-g43ff75f2c3f (aka 11.0.1 rc2). --- clang/include/clang/AST/DeclTemplate.h | 0 clang/include/clang/Basic/Attr.td | 10 +- clang/include/clang/Basic/AttrDocs.td | 26 +++ clang/include/clang/Basic/Builtins.def | 27 ++- clang/include/clang/Basic/Builtins.h | 7 + clang/include/clang/Basic/DiagnosticDriverKinds.td | 2 +- clang/include/clang/Basic/IdentifierTable.h | 12 -- clang/include/clang/Format/Format.h | 0 clang/include/clang/Sema/Sema.h | 2 + .../Tooling/Refactoring/Extract/SourceExtraction.h | 0 clang/lib/AST/Decl.cpp | 28 +-- clang/lib/AST/DeclTemplate.cpp | 0 clang/lib/CodeGen/CGBuiltin.cpp | 6 +- clang/lib/CodeGen/CGObjCGNU.cpp | 13 -- clang/lib/CodeGen/TargetInfo.cpp | 16 +- clang/lib/Driver/ToolChains/Cuda.cpp | 113 ++++++++--- clang/lib/Driver/ToolChains/Cuda.h | 3 - clang/lib/Format/FormatToken.h | 6 - clang/lib/Format/UnwrappedLineParser.cpp | 2 - clang/lib/Format/WhitespaceManager.cpp | 10 +- clang/lib/Headers/__clang_cuda_complex_builtins.h | 6 +- clang/lib/Headers/cuda_wrappers/new | 8 + clang/lib/Headers/intrin.h | 145 ++++++-------- clang/lib/Headers/openmp_wrappers/complex | 2 + clang/lib/Headers/openmp_wrappers/complex.h | 2 + clang/lib/Sema/SemaChecking.cpp | 5 - clang/lib/Sema/SemaConcept.cpp | 0 clang/lib/Sema/SemaDecl.cpp | 129 +++++++----- clang/lib/Sema/SemaDeclAttr.cpp | 6 +- clang/lib/Sema/SemaExpr.cpp | 3 +- clang/lib/Sema/SemaLookup.cpp | 7 +- clang/lib/Sema/SemaTemplate.cpp | 3 +- clang/lib/Sema/SemaTemplateInstantiateDecl.cpp | 13 +- clang/lib/Sema/UsedDeclVisitor.h | 11 +- clang/lib/Serialization/ASTReader.cpp | 12 +- clang/lib/Serialization/ASTWriter.cpp | 6 +- clang/lib/StaticAnalyzer/Checkers/Yaml.h | 0 .../sanitizer_common_interceptors_ioctl.inc | 0 lld/COFF/DebugTypes.cpp | 5 +- lld/ELF/LinkerScript.cpp | 7 +- lldb/include/lldb/Core/ModuleList.h | 19 +- lldb/include/lldb/Target/Platform.h | 9 +- lldb/source/Core/ModuleList.cpp | 39 ++-- .../Plugins/Process/minidump/ProcessMinidump.cpp | 152 +++++++++++--- .../Plugins/Process/minidump/ProcessMinidump.h | 4 + lldb/source/Target/Platform.cpp | 43 ++-- lldb/source/Target/Target.cpp | 98 +++++++-- llvm/include/llvm/ADT/ImmutableMap.h | 2 +- llvm/include/llvm/BinaryFormat/ELF.h | 6 + .../llvm/Transforms/IPO/DeadArgumentElimination.h | 1 + llvm/lib/DebugInfo/DWARF/DWARFDebugArangeSet.cpp | 13 +- llvm/lib/DebugInfo/DWARF/DWARFUnitIndex.cpp | 8 +- .../ExecutionEngine/RuntimeDyld/RuntimeDyld.cpp | 19 +- llvm/lib/IR/ConstantFold.cpp | 2 +- llvm/lib/MC/MCObjectFileInfo.cpp | 2 + llvm/lib/MC/MCParser/ELFAsmParser.cpp | 7 +- llvm/lib/MCA/HardwareUnits/LSUnit.cpp | 2 + llvm/lib/Support/Host.cpp | 5 +- llvm/lib/Support/Windows/Path.inc | 21 +- llvm/lib/Target/AArch64/AArch64AsmPrinter.cpp | 15 ++ .../Target/AArch64/AArch64CallingConvention.cpp | 9 + llvm/lib/Target/AArch64/AArch64ISelLowering.cpp | 17 +- .../Target/AArch64/AsmParser/AArch64AsmParser.cpp | 30 +++ .../AArch64/GISel/AArch64InstructionSelector.cpp | 134 +++++-------- .../Target/AArch64/GISel/AArch64LegalizerInfo.cpp | 38 ++-- .../AArch64/GISel/AArch64RegisterBankInfo.cpp | 1 + .../AArch64/MCTargetDesc/AArch64ELFStreamer.cpp | 8 + .../AArch64/MCTargetDesc/AArch64TargetStreamer.h | 4 + llvm/lib/Target/AMDGPU/SIPreEmitPeephole.cpp | 24 ++- llvm/lib/Target/BPF/BPFISelDAGToDAG.cpp | 4 +- llvm/lib/Target/Hexagon/HexagonAsmPrinter.h | 0 llvm/lib/Target/Sparc/LeonFeatures.td | 0 llvm/lib/Target/Sparc/LeonPasses.cpp | 0 llvm/lib/Target/Sparc/LeonPasses.h | 0 llvm/lib/Target/Sparc/SparcAsmPrinter.cpp | 2 +- llvm/lib/Target/Sparc/SparcSchedule.td | 0 .../lib/Target/WebAssembly/WebAssemblyFastISel.cpp | 17 +- .../lib/Target/WebAssembly/WebAssemblyInstrInfo.td | 4 +- .../Target/WebAssembly/WebAssemblyRegisterInfo.cpp | 10 +- llvm/lib/Target/X86/X86EvexToVex.cpp | 0 llvm/lib/Target/X86/X86FrameLowering.cpp | 222 +++++++++++++++++++-- llvm/lib/Target/X86/X86FrameLowering.h | 8 +- llvm/lib/Target/X86/X86ISelLowering.cpp | 29 ++- llvm/lib/Target/X86/X86SchedBroadwell.td | 0 llvm/lib/Target/X86/X86SchedSkylakeServer.td | 0 .../lib/Transforms/IPO/DeadArgumentElimination.cpp | 29 ++- .../Transforms/InstCombine/InstCombineSelect.cpp | 20 +- .../Instrumentation/AddressSanitizer.cpp | 20 +- 88 files changed, 1152 insertions(+), 598 deletions(-) diff --git a/clang/include/clang/AST/DeclTemplate.h b/clang/include/clang/AST/DeclTemplate.h old mode 100644 new mode 100755 diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index bc4a380545af..60eaee7839e2 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -2115,7 +2115,7 @@ def InitPriority : InheritableAttr { let Spellings = [GCC<"init_priority", /*AllowInC*/0>]; let Args = [UnsignedArgument<"Priority">]; let Subjects = SubjectList<[Var], ErrorDiag>; - let Documentation = [Undocumented]; + let Documentation = [InitPriorityDocs]; } def Section : InheritableAttr { @@ -3444,3 +3444,11 @@ def ReleaseHandle : InheritableParamAttr { let Subjects = SubjectList<[ParmVar]>; let Documentation = [ReleaseHandleDocs]; } + +def Builtin : InheritableAttr { + let Spellings = []; + let Args = [UnsignedArgument<"ID">]; + let Subjects = SubjectList<[Function]>; + let SemaHandler = 0; + let Documentation = [Undocumented]; +} diff --git a/clang/include/clang/Basic/AttrDocs.td b/clang/include/clang/Basic/AttrDocs.td index 3cba3a3d96f9..833127ed44eb 100644 --- a/clang/include/clang/Basic/AttrDocs.td +++ b/clang/include/clang/Basic/AttrDocs.td @@ -57,6 +57,32 @@ global variable or function should be in after translation. let Heading = "section, __declspec(allocate)"; } +def InitPriorityDocs : Documentation { + let Category = DocCatVariable; + let Content = [{ +In C++, the order in which global variables are initialized across translation +units is unspecified, unlike the ordering within a single translation unit. The +``init_priority`` attribute allows you to specify a relative ordering for the +initialization of objects declared at namespace scope in C++. The priority is +given as an integer constant expression between 101 and 65535 (inclusive). +Priorities outside of that range are reserved for use by the implementation. A +lower value indicates a higher priority of initialization. Note that only the +relative ordering of values is important. For example: + +.. code-block:: c++ + + struct SomeType { SomeType(); }; + __attribute__((init_priority(200))) SomeType Obj1; + __attribute__((init_priority(101))) SomeType Obj2; + +``Obj1`` will be initialized *before* ``Obj2`` despite the usual order of +initialization being the opposite. + +This attribute is only supported for C++ and Objective-C++ and is ignored in +other language modes. + }]; +} + def InitSegDocs : Documentation { let Category = DocCatVariable; let Content = [{ diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def index 1416a64543a4..5463b7dfc18c 100644 --- a/clang/include/clang/Basic/Builtins.def +++ b/clang/include/clang/Basic/Builtins.def @@ -75,6 +75,9 @@ // U -> pure // c -> const // t -> signature is meaningless, use custom typechecking +// T -> type is not important to semantic analysis and codegen; recognize as +// builtin even if type doesn't match signature, and don't warn if we +// can't be sure the type is right // F -> this is a libc/libm function with a '__builtin_' prefix added. // f -> this is a libc/libm function without the '__builtin_' prefix. It can // be followed by ':headername:' to state which header this function @@ -893,7 +896,7 @@ LANGBUILTIN(__va_start, "vc**.", "nt", ALL_MS_LANGUAGES) LANGBUILTIN(__fastfail, "vUi", "nr", ALL_MS_LANGUAGES) // Microsoft library builtins. -LIBBUILTIN(_setjmpex, "iJ", "fj", "setjmpex.h", ALL_MS_LANGUAGES) +LIBBUILTIN(_setjmpex, "iJ", "fjT", "setjmpex.h", ALL_MS_LANGUAGES) // C99 library functions // C99 stdarg.h @@ -987,8 +990,8 @@ LIBBUILTIN(wmemmove,"w*w*wC*z", "f", "wchar.h", ALL_LANGUAGES) // In some systems setjmp is a macro that expands to _setjmp. We undefine // it here to avoid having two identical LIBBUILTIN entries. #undef setjmp -LIBBUILTIN(setjmp, "iJ", "fj", "setjmp.h", ALL_LANGUAGES) -LIBBUILTIN(longjmp, "vJi", "fr", "setjmp.h", ALL_LANGUAGES) +LIBBUILTIN(setjmp, "iJ", "fjT", "setjmp.h", ALL_LANGUAGES) +LIBBUILTIN(longjmp, "vJi", "frT", "setjmp.h", ALL_LANGUAGES) // Non-C library functions, active in GNU mode only. // Functions with (returns_twice) attribute (marked as "j") are still active in @@ -1015,20 +1018,22 @@ LIBBUILTIN(strcasecmp, "icC*cC*", "f", "strings.h", ALL_GNU_LANGUAGES) LIBBUILTIN(strncasecmp, "icC*cC*z", "f", "strings.h", ALL_GNU_LANGUAGES) // POSIX unistd.h LIBBUILTIN(_exit, "vi", "fr", "unistd.h", ALL_GNU_LANGUAGES) -LIBBUILTIN(vfork, "p", "fj", "unistd.h", ALL_LANGUAGES) +LIBBUILTIN(vfork, "p", "fjT", "unistd.h", ALL_LANGUAGES) // POSIX pthread.h +// FIXME: Should specify argument types. LIBBUILTIN(pthread_create, "", "fC<2,3>", "pthread.h", ALL_GNU_LANGUAGES) // POSIX setjmp.h -LIBBUILTIN(_setjmp, "iJ", "fj", "setjmp.h", ALL_LANGUAGES) -LIBBUILTIN(__sigsetjmp, "iSJi", "fj", "setjmp.h", ALL_LANGUAGES) -LIBBUILTIN(sigsetjmp, "iSJi", "fj", "setjmp.h", ALL_LANGUAGES) -LIBBUILTIN(savectx, "iJ", "fj", "setjmp.h", ALL_LANGUAGES) -LIBBUILTIN(getcontext, "iK*", "fj", "setjmp.h", ALL_LANGUAGES) +// FIXME: MinGW _setjmp has an additional void* parameter. +LIBBUILTIN(_setjmp, "iJ", "fjT", "setjmp.h", ALL_LANGUAGES) +LIBBUILTIN(__sigsetjmp, "iSJi", "fjT", "setjmp.h", ALL_LANGUAGES) +LIBBUILTIN(sigsetjmp, "iSJi", "fjT", "setjmp.h", ALL_LANGUAGES) +LIBBUILTIN(savectx, "iJ", "fjT", "setjmp.h", ALL_LANGUAGES) +LIBBUILTIN(getcontext, "iK*", "fjT", "setjmp.h", ALL_LANGUAGES) -LIBBUILTIN(_longjmp, "vJi", "fr", "setjmp.h", ALL_GNU_LANGUAGES) -LIBBUILTIN(siglongjmp, "vSJi", "fr", "setjmp.h", ALL_GNU_LANGUAGES) +LIBBUILTIN(_longjmp, "vJi", "frT", "setjmp.h", ALL_GNU_LANGUAGES) +LIBBUILTIN(siglongjmp, "vSJi", "frT", "setjmp.h", ALL_GNU_LANGUAGES) // non-standard but very common LIBBUILTIN(strlcpy, "zc*cC*z", "f", "string.h", ALL_GNU_LANGUAGES) LIBBUILTIN(strlcat, "zc*cC*z", "f", "string.h", ALL_GNU_LANGUAGES) diff --git a/clang/include/clang/Basic/Builtins.h b/clang/include/clang/Basic/Builtins.h index e4ed482d9068..15bfcf797917 100644 --- a/clang/include/clang/Basic/Builtins.h +++ b/clang/include/clang/Basic/Builtins.h @@ -158,6 +158,13 @@ public: return strchr(getRecord(ID).Attributes, 't') != nullptr; } + /// Determines whether a declaration of this builtin should be recognized + /// even if the type doesn't match the specified signature. + bool allowTypeMismatch(unsigned ID) const { + return strchr(getRecord(ID).Attributes, 'T') != nullptr || + hasCustomTypechecking(ID); + } + /// Determines whether this builtin has a result or any arguments which /// are pointer types. bool hasPtrArgsOrResult(unsigned ID) const { diff --git a/clang/include/clang/Basic/DiagnosticDriverKinds.td b/clang/include/clang/Basic/DiagnosticDriverKinds.td index 558639ecad6a..acdad15cdf6c 100644 --- a/clang/include/clang/Basic/DiagnosticDriverKinds.td +++ b/clang/include/clang/Basic/DiagnosticDriverKinds.td @@ -69,7 +69,7 @@ def err_drv_cuda_version_unsupported : Error< "install, pass a different GPU arch with --cuda-gpu-arch, or pass " "--no-cuda-version-check.">; def warn_drv_unknown_cuda_version: Warning< - "Unknown CUDA version %0. Assuming the latest supported version %1">, + "Unknown CUDA version. %0 Assuming the latest supported version %1">, InGroup<CudaUnknownVersion>; def err_drv_cuda_host_arch : Error<"unsupported architecture '%0' for host compilation.">; def err_drv_mix_cuda_hip : Error<"Mixed Cuda and HIP compilation is not supported.">; diff --git a/clang/include/clang/Basic/IdentifierTable.h b/clang/include/clang/Basic/IdentifierTable.h index fc554a35e721..204a0f0cc0a5 100644 --- a/clang/include/clang/Basic/IdentifierTable.h +++ b/clang/include/clang/Basic/IdentifierTable.h @@ -225,18 +225,6 @@ public: } void setObjCKeywordID(tok::ObjCKeywordKind ID) { ObjCOrBuiltinID = ID; } - /// True if setNotBuiltin() was called. - bool hasRevertedBuiltin() const { - return ObjCOrBuiltinID == tok::NUM_OBJC_KEYWORDS; - } - - /// Revert the identifier to a non-builtin identifier. We do this if - /// the name of a known builtin library function is used to declare that - /// function, but an unexpected type is specified. - void revertBuiltin() { - setBuiltinID(0); - } - /// Return a value indicating whether this is a builtin function. /// /// 0 is not-built-in. 1+ are specific builtin functions. diff --git a/clang/include/clang/Format/Format.h b/clang/include/clang/Format/Format.h old mode 100644 new mode 100755 diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h index 6f7ad8076718..16a7084f6b08 100644 --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -3957,6 +3957,8 @@ public: ObjCInterfaceDecl *getObjCInterfaceDecl(IdentifierInfo *&Id, SourceLocation IdLoc, bool TypoCorrection = false); + FunctionDecl *CreateBuiltin(IdentifierInfo *II, QualType Type, unsigned ID, + SourceLocation Loc); NamedDecl *LazilyCreateBuiltin(IdentifierInfo *II, unsigned ID, Scope *S, bool ForRedeclaration, SourceLocation Loc); diff --git a/clang/include/clang/Tooling/Refactoring/Extract/SourceExtraction.h b/clang/include/clang/Tooling/Refactoring/Extract/SourceExtraction.h old mode 100755 new mode 100644 diff --git a/clang/lib/AST/Decl.cpp b/clang/lib/AST/Decl.cpp index 5c0a98815dd7..67490da5c584 100644 --- a/clang/lib/AST/Decl.cpp +++ b/clang/lib/AST/Decl.cpp @@ -3162,44 +3162,24 @@ FunctionDecl *FunctionDecl::getCanonicalDecl() { return getFirstDecl(); } /// functions as their wrapped builtins. This shouldn't be done in general, but /// it's useful in Sema to diagnose calls to wrappers based on their semantics. unsigned FunctionDecl::getBuiltinID(bool ConsiderWrapperFunctions) const { - unsigned BuiltinID; + unsigned BuiltinID = 0; if (const auto *ABAA = getAttr<ArmBuiltinAliasAttr>()) { BuiltinID = ABAA->getBuiltinName()->getBuiltinID(); - } else { - if (!getIdentifier()) - return 0; - - BuiltinID = getIdentifier()->getBuiltinID(); + } else if (const auto *A = getAttr<BuiltinAttr>()) { + BuiltinID = A->getID(); } if (!BuiltinID) return 0; - ASTContext &Context = getASTContext(); - if (Context.getLangOpts().CPlusPlus) { - const auto *LinkageDecl = - dyn_cast<LinkageSpecDecl>(getFirstDecl()->getDeclContext()); - // In C++, the first declaration of a builtin is always inside an implicit - // extern "C". - // FIXME: A recognised library function may not be directly in an extern "C" - // declaration, for instance "extern "C" { namespace std { decl } }". - if (!LinkageDecl) { - if (BuiltinID == Builtin::BI__GetExceptionInfo && - Context.getTargetInfo().getCXXABI().isMicrosoft()) - return Builtin::BI__GetExceptionInfo; - return 0; - } - if (LinkageDecl->getLanguage() != LinkageSpecDecl::lang_c) - return 0; - } - // If the function is marked "overloadable", it has a different mangled name // and is not the C library function. if (!ConsiderWrapperFunctions && hasAttr<OverloadableAttr>() && !hasAttr<ArmBuiltinAliasAttr>()) return 0; + ASTContext &Context = getASTContext(); if (!Context.BuiltinInfo.isPredefinedLibFunction(BuiltinID)) return BuiltinID; diff --git a/clang/lib/AST/DeclTemplate.cpp b/clang/lib/AST/DeclTemplate.cpp old mode 100644 new mode 100755 diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 8994b939093e..6b93f1b60af5 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -3754,11 +3754,13 @@ RValue CodeGenFunction::EmitBuiltinExpr(const GlobalDecl GD, unsigned BuiltinID, case Builtin::BI_abnormal_termination: return RValue::get(EmitSEHAbnormalTermination()); case Builtin::BI_setjmpex: - if (getTarget().getTriple().isOSMSVCRT()) + if (getTarget().getTriple().isOSMSVCRT() && E->getNumArgs() == 1 && + E->getArg(0)->getType()->isPointerType()) return EmitMSVCRTSetJmp(*this, MSVCSetJmpKind::_setjmpex, E); break; case Builtin::BI_setjmp: - if (getTarget().getTriple().isOSMSVCRT()) { + if (getTarget().getTriple().isOSMSVCRT() && E->getNumArgs() == 1 && + E->getArg(0)->getType()->isPointerType()) { if (getTarget().getTriple().getArch() == llvm::Triple::x86) return EmitMSVCRTSetJmp(*this, MSVCSetJmpKind::_setjmp3, E); else if (getTarget().getTriple().getArch() == llvm::Triple::aarch64) diff --git a/clang/lib/CodeGen/CGObjCGNU.cpp b/clang/lib/CodeGen/CGObjCGNU.cpp index bb9c494ae68e..c64faf4c0af7 100644 --- a/clang/lib/CodeGen/CGObjCGNU.cpp +++ b/clang/lib/CodeGen/CGObjCGNU.cpp @@ -3511,19 +3511,6 @@ void CGObjCGNU::GenerateClass(const ObjCImplementationDecl *OID) { ClassMethods.insert(ClassMethods.begin(), OID->classmeth_begin(), OID->classmeth_end()); - // Collect the same information about synthesized properties, which don't - // show up in the instance method lists. - for (auto *propertyImpl : OID->property_impls()) - if (propertyImpl->getPropertyImplementation() == - ObjCPropertyImplDecl::Synthesize) { - auto addPropertyMethod = [&](const ObjCMethodDecl *accessor) { - if (accessor) - InstanceMethods.push_back(accessor); - }; - addPropertyMethod(propertyImpl->getGetterMethodDecl()); - addPropertyMethod(propertyImpl->getSetterMethodDecl()); - } - llvm::Constant *Properties = GeneratePropertyList(OID, ClassDecl); // Collect the names of referenced protocols diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp index 9cd63ebe29ee..a061651d8b21 100644 --- a/clang/lib/CodeGen/TargetInfo.cpp +++ b/clang/lib/CodeGen/TargetInfo.cpp @@ -5471,6 +5471,11 @@ private: Address EmitVAArg(CodeGenFunction &CGF, Address VAListAddr, QualType Ty) const override { + llvm::Type *BaseTy = CGF.ConvertType(Ty); + if (isa<llvm::ScalableVectorType>(BaseTy)) + llvm::report_fatal_error("Passing SVE types to variadic functions is " + "currently not supported"); + return Kind == Win64 ? EmitMSVAArg(CGF, VAListAddr, Ty) : isDarwinPCS() ? EmitDarwinVAArg(VAListAddr, Ty, CGF) : EmitAAPCSVAArg(VAListAddr, Ty, CGF); @@ -10390,7 +10395,6 @@ bool RISCVABIInfo::detectFPCCEligibleStructHelper(QualType Ty, CharUnits CurOff, return false; Field1Ty = CGT.ConvertType(EltTy); Field1Off = CurOff; - assert(CurOff.isZero() && "Unexpected offset for first field"); Field2Ty = Field1Ty; Field2Off = Field1Off + getContext().getTypeSizeInChars(EltTy); return true; @@ -10485,7 +10489,7 @@ bool RISCVABIInfo::detectFPCCEligibleStruct(QualType Ty, llvm::Type *&Field1Ty, NeededArgFPRs++; else if (Field2Ty) NeededArgGPRs++; - return IsCandidate; + return true; } // Call getCoerceAndExpand for the two-element flattened struct described by @@ -10511,15 +10515,15 @@ ABIArgInfo RISCVABIInfo::coerceAndExpandFPCCEligibleStruct( CharUnits Field2Align = CharUnits::fromQuantity(getDataLayout().getABITypeAlignment(Field2Ty)); - CharUnits Field1Size = + CharUnits Field1End = Field1Off + CharUnits::fromQuantity(getDataLayout().getTypeStoreSize(Field1Ty)); - CharUnits Field2OffNoPadNoPack = Field1Size.alignTo(Field2Align); + CharUnits Field2OffNoPadNoPack = Field1End.alignTo(Field2Align); CharUnits Padding = CharUnits::Zero(); if (Field2Off > Field2OffNoPadNoPack) Padding = Field2Off - Field2OffNoPadNoPack; - else if (Field2Off != Field2Align && Field2Off > Field1Size) - Padding = Field2Off - Field1Size; + else if (Field2Off != Field2Align && Field2Off > Field1End) + Padding = Field2Off - Field1End; bool IsPacked = !Field2Off.isMultipleOf(Field2Align); diff --git a/clang/lib/Driver/ToolChains/Cuda.cpp b/clang/lib/Driver/ToolChains/Cuda.cpp index 110a0bca9bc1..ffc606dd554b 100644 --- a/clang/lib/Driver/ToolChains/Cuda.cpp +++ b/clang/lib/Driver/ToolChains/Cuda.cpp @@ -16,6 +16,7 @@ #include "clang/Driver/Driver.h" #include "clang/Driver/DriverDiagnostic.h" #include "clang/Driver/Options.h" +#include "llvm/ADT/Optional.h" #include "llvm/Option/ArgList.h" #include "llvm/Support/FileSystem.h" #include "llvm/Support/Host.h" @@ -32,30 +33,81 @@ using namespace clang::driver::tools; using namespace clang; using namespace llvm::opt; +namespace { +struct CudaVersionInfo { + std::string DetectedVersion; + CudaVersion Version; +}; // Parses the contents of version.txt in an CUDA installation. It should // contain one line of the from e.g. "CUDA Version 7.5.2". -void CudaInstallationDetector::ParseCudaVersionFile(llvm::StringRef V) { - Version = CudaVersion::UNKNOWN; +CudaVersionInfo parseCudaVersionFile(llvm::StringRef V) { + V = V.trim(); if (!V.startswith("CUDA Version ")) - return; + return {V.str(), CudaVersion::UNKNOWN}; V = V.substr(strlen("CUDA Version ")); SmallVector<StringRef,4> VersionParts; V.split(VersionParts, '.'); - if (VersionParts.size() < 2) - return; - DetectedVersion = join_items(".", VersionParts[0], VersionParts[1]); - Version = CudaStringToVersion(DetectedVersion); - if (Version != CudaVersion::UNKNOWN) { - // TODO(tra): remove the warning once we have all features of 10.2 and 11.0 - // implemented. - DetectedVersionIsNotSupported = Version > CudaVersion::LATEST_SUPPORTED; - return; - } + return {"version.txt: " + V.str() + ".", + VersionParts.size() < 2 + ? CudaVersion::UNKNOWN + : CudaStringToVersion( + join_items(".", VersionParts[0], VersionParts[1]))}; +} - Version = CudaVersion::LATEST_SUPPORTED; - DetectedVersionIsNotSupported = true; +CudaVersion getCudaVersion(uint32_t raw_version) { + if (raw_version < 7050) + return CudaVersion::CUDA_70; + if (raw_version < 8000) + return CudaVersion::CUDA_75; + if (raw_version < 9000) + return CudaVersion::CUDA_80; + if (raw_version < 9010) + return CudaVersion::CUDA_90; + if (raw_version < 9020) + return CudaVersion::CUDA_91; + if (raw_version < 10000) + return CudaVersion::CUDA_92; + if (raw_version < 10010) + return CudaVersion::CUDA_100; + if (raw_version < 10020) + return CudaVersion::CUDA_101; + if (raw_version < 11000) + return CudaVersion::CUDA_102; + if (raw_version < 11010) + return CudaVersion::CUDA_110; + return CudaVersion::LATEST; } +CudaVersionInfo parseCudaHFile(llvm::StringRef Input) { + // Helper lambda which skips the words if the line starts with them or returns + // None otherwise. + auto StartsWithWords = + [](llvm::StringRef Line, + const SmallVector<StringRef, 3> words) -> llvm::Optional<StringRef> { + for (StringRef word : words) { + if (!Line.consume_front(word)) + return {}; + Line = Line.ltrim(); + } + return Line; + }; + + Input = Input.ltrim(); + while (!Input.empty()) { + if (auto Line = + StartsWithWords(Input.ltrim(), {"#", "define", "CUDA_VERSION"})) { + uint32_t RawVersion; + Line->consumeInteger(10, RawVersion); + return {"cuda.h: CUDA_VERSION=" + Twine(RawVersion).str() + ".", + getCudaVersion(RawVersion)}; + } + // Find next non-empty line. + Input = Input.drop_front(Input.find_first_of("\n\r")).ltrim(); + } + return {"cuda.h: CUDA_VERSION not found.", CudaVersion::UNKNOWN}; +} +} // namespace + void CudaInstallationDetector::WarnIfUnsupportedVersion() { if (DetectedVersionIsNotSupported) D.Diag(diag::warn_drv_unknown_cuda_version) @@ -152,16 +204,31 @@ CudaInstallationDetector::CudaInstallationDetector( else continue; - llvm::ErrorOr<std::unique_ptr<llvm::MemoryBuffer>> VersionFile = - FS.getBufferForFile(InstallPath + "/version.txt"); - if (!VersionFile) { - // CUDA 7.0 doesn't have a version.txt, so guess that's our version if - // version.txt isn't present. - Version = CudaVersion::CUDA_70; - } else { - ParseCudaVersionFile((*VersionFile)->getBuffer()); + CudaVersionInfo VersionInfo = {"", CudaVersion::UNKNOWN}; + if (auto VersionFile = FS.getBufferForFile(InstallPath + "/version.txt")) + VersionInfo = parseCudaVersionFile((*VersionFile)->getBuffer()); + // If version file didn't give us the version, try to find it in cuda.h + if (VersionInfo.Version == CudaVersion::UNKNOWN) + if (auto CudaHFile = FS.getBufferForFile(InstallPath + "/include/cuda.h")) + VersionInfo = parseCudaHFile((*CudaHFile)->getBuffer()); + // As the last resort, make an educated guess between CUDA-7.0, (which had + // no version.txt file and had old-style libdevice bitcode ) and an unknown + // recent CUDA version (no version.txt, new style bitcode). + if (VersionInfo.Version == CudaVersion::UNKNOWN) { + VersionInfo.Version = (FS.exists(LibDevicePath + "/libdevice.10.bc")) + ? Version = CudaVersion::LATEST + : Version = CudaVersion::CUDA_70; + VersionInfo.DetectedVersion = + "No version found in version.txt or cuda.h."; } + Version = VersionInfo.Version; + DetectedVersion = VersionInfo.DetectedVersion; + + // TODO(tra): remove the warning once we have all features of 10.2 + // and 11.0 implemented. + DetectedVersionIsNotSupported = Version > CudaVersion::LATEST_SUPPORTED; + if (Version >= CudaVersion::CUDA_90) { // CUDA-9+ uses single libdevice file for all GPU variants. std::string FilePath = LibDevicePath + "/libdevice.10.bc"; diff --git a/clang/lib/Driver/ToolChains/Cuda.h b/clang/lib/Driver/ToolChains/Cuda.h index 873eb7338a30..bbf272c468a5 100644 --- a/clang/lib/Driver/ToolChains/Cuda.h +++ b/clang/lib/Driver/ToolChains/Cuda.h @@ -78,9 +78,6 @@ public: return LibDeviceMap.lookup(Gpu); } void WarnIfUnsupportedVersion(); - -private: - void ParseCudaVersionFile(llvm::StringRef V); }; namespace tools { diff --git a/clang/lib/Format/FormatToken.h b/clang/lib/Format/FormatToken.h index d4287f53fde3..b132a3e84da5 100644 --- a/clang/lib/Format/FormatToken.h +++ b/clang/lib/Format/FormatToken.h @@ -183,12 +183,6 @@ struct FormatToken { /// before the token. bool MustBreakBefore = false; - /// Whether to not align across this token - /// - /// This happens for example when a preprocessor directive ended directly - /// before the token, but very rarely otherwise. - bool MustBreakAlignBefore = false; - /// The raw text of the token. /// /// Contains the raw token text without leading whitespace and without leading diff --git a/clang/lib/Format/UnwrappedLineParser.cpp b/clang/lib/Format/UnwrappedLineParser.cpp index a37386425aae..ea8a41cfba82 100644 --- a/clang/lib/Format/UnwrappedLineParser.cpp +++ b/clang/lib/Format/UnwrappedLineParser.cpp @@ -3037,7 +3037,6 @@ void UnwrappedLineParser::readToken(int LevelDifference) { } FormatTok = Tokens->getNextToken(); FormatTok->MustBreakBefore = true; - FormatTok->MustBreakAlignBefore = true; } if (!PPStack.empty() && (PPStack.back().Kind == PP_Unreachable) && @@ -3062,7 +3061,6 @@ void UnwrappedLineParser::pushToken(FormatToken *Tok) { Line->Tokens.push_back(UnwrappedLineNode(Tok)); if (MustBreakBeforeNextToken) { Line->Tokens.back().Tok->MustBreakBefore = true; - Line->Tokens.back().Tok->MustBreakAlignBefore = true; MustBreakBeforeNextToken = false; } } diff --git a/clang/lib/Format/WhitespaceManager.cpp b/clang/lib/Format/WhitespaceManager.cpp index 32e0b685ea0f..3a265bd09168 100644 --- a/clang/lib/Format/WhitespaceManager.cpp +++ b/clang/lib/Format/WhitespaceManager.cpp @@ -411,11 +411,9 @@ static unsigned AlignTokens(const FormatStyle &Style, F &&Matches, if (Changes[i].NewlinesBefore != 0) { CommasBeforeMatch = 0; EndOfSequence = i; - // If there is a blank line, there is a forced-align-break (eg, - // preprocessor), or if the last line didn't contain any matching token, - // the sequence ends here. - if (Changes[i].NewlinesBefore > 1 || - Changes[i].Tok->MustBreakAlignBefore || !FoundMatchOnLine) + // If there is a blank line, or if the last line didn't contain any + // matching token, the sequence ends here. + if (Changes[i].NewlinesBefore > 1 || !FoundMatchOnLine) AlignCurrentSequence(); FoundMatchOnLine = false; @@ -726,8 +724,6 @@ void WhitespaceManager::alignTrailingComments() { if (Changes[i].StartOfBlockComment) continue; Newlines += Changes[i].NewlinesBefore; - if (Changes[i].Tok->MustBreakAlignBefore) - BreakBeforeNext = true; if (!Changes[i].IsTrailingComment) continue; diff --git a/clang/lib/Headers/__clang_cuda_complex_builtins.h b/clang/lib/Headers/__clang_cuda_complex_builtins.h index 8c10ff6b461f..d924487ab285 100644 --- a/clang/lib/Headers/__clang_cuda_complex_builtins.h +++ b/clang/lib/Headers/__clang_cuda_complex_builtins.h @@ -16,7 +16,7 @@ // to work with CUDA and OpenMP target offloading [in C and C++ mode].) #pragma push_macro("__DEVICE__") -#ifdef _OPENMP +#ifdef __OPENMP_NVPTX__ #pragma omp declare target #define __DEVICE__ __attribute__((noinline, nothrow, cold, weak)) #else @@ -26,7 +26,7 @@ // To make the algorithms available for C and C++ in CUDA and OpenMP we select // different but equivalent function versions. TODO: For OpenMP we currently // select the native builtins as the overload support for templates is lacking. -#if !defined(_OPENMP) +#if !defined(__OPENMP_NVPTX__) #define _ISNANd std::isnan #define _ISNANf std::isnan #define _ISINFd std::isinf @@ -250,7 +250,7 @@ __DEVICE__ float _Complex __divsc3(float __a, float __b, float __c, float __d) { #undef _LOGBd #undef _LOGBf -#ifdef _OPENMP +#ifdef __OPENMP_NVPTX__ #pragma omp end declare target #endif diff --git a/clang/lib/Headers/cuda_wrappers/new b/clang/lib/Headers/cuda_wrappers/new index f49811c5a57c..7f255314056a 100644 --- a/clang/lib/Headers/cuda_wrappers/new +++ b/clang/lib/Headers/cuda_wrappers/new @@ -26,6 +26,13 @@ #include_next <new> +#if !defined(__device__) +// The header has been included too early from the standard C++ library +// and CUDA-specific macros are not available yet. +// Undo the include guard and try again later. +#undef __CLANG_CUDA_WRAPPERS_NEW +#else + #pragma push_macro("CUDA_NOEXCEPT") #if __cplusplus >= 201103L #define CUDA_NOEXCEPT noexcept @@ -95,4 +102,5 @@ __device__ inline void operator delete[](void *, void *) CUDA_NOEXCEPT {} #pragma pop_macro("CUDA_NOEXCEPT") +#endif // __device__ #endif // include guard diff --git a/clang/lib/Headers/intrin.h b/clang/lib/Headers/intrin.h index 871b47ca8267..e7b76a3bb2ed 100644 --- a/clang/lib/Headers/intrin.h +++ b/clang/lib/Headers/intrin.h @@ -57,16 +57,11 @@ void __addfsbyte(unsigned long, unsigned char); void __addfsdword(unsigned long, unsigned long); void __addfsword(unsigned long, unsigned short); void __code_seg(const char *); -static __inline__ void __cpuid(int[4], int); -static __inline__ void __cpuidex(int[4], int, int); -static __inline__ __int64 __emul(int, int); -static __inline__ unsigned __int64 __emulu(unsigned int, unsigned int); unsigned int __getcallerseflags(void); -static __inline__ void __halt(void); unsigned char __inbyte(unsigned short); void __inbytestring(unsigned short, unsigned char *, unsigned long); @@ -82,13 +77,9 @@ void __inwordstring(unsigned short, unsigned short *, unsigned long); void __lidt(void *); unsigned __int64 __ll_lshift(unsigned __int64, int); __int64 __ll_rshift(__int64, int); -static __inline__ void __movsb(unsigned char *, unsigned char const *, size_t); -static __inline__ void __movsd(unsigned long *, unsigned long const *, size_t); -static __inline__ void __movsw(unsigned short *, unsigned short const *, size_t); -static __inline__ void __nop(void); void __nvreg_restore_fence(void); void __nvreg_save_fence(void); @@ -105,23 +96,16 @@ unsigned long __readcr4(void); unsigned long __readcr8(void); unsigned int __readdr(unsigned int); #ifdef __i386__ -static __inline__ unsigned char __readfsbyte(unsigned long); -static __inline__ unsigned __int64 __readfsqword(unsigned long); -static __inline__ unsigned short __readfsword(unsigned long); #endif -static __inline__ unsigned __int64 __readmsr(unsigned long); unsigned __int64 __readpmc(unsigned long); unsigned long __segmentlimit(unsigned long); void __sidt(void *); -static __inline__ void __stosb(unsigned char *, unsigned char, size_t); -static __inline__ void __stosd(unsigned long *, unsigned long, size_t); -static __inline__ void __stosw(unsigned short *, unsigned short, size_t); void __svm_clgi(void); void __svm_invlpga(void *, int); @@ -136,7 +120,6 @@ void __vmx_off(void); void __vmx_vmptrst(unsigned __int64 *); void __wbinvd(void); void __writecr0(unsigned int); -static __inline__ void __writecr3(unsigned __INTPTR_TYPE__); void __writecr4(unsigned int); void __writecr8(unsigned int); @@ -146,11 +129,8 @@ void __writefsdword(unsigned long, unsigned long); void __writefsqword(unsigned long, unsigned __int64); void __writefsword(unsigned long, unsigned short); void __writemsr(unsigned long, unsigned __int64); -static __inline__ void *_AddressOfReturnAddress(void); -static __inline__ unsigned char _BitScanForward(unsigned long *_Index, unsigned long _Mask); -static __inline__ unsigned char _BitScanReverse(unsigned long *_Index, unsigned long _Mask); unsigned char _bittest(long const *, long); unsigned char _bittestandcomplement(long *, long); @@ -169,12 +149,10 @@ long _InterlockedExchangeAdd_HLEAcquire(long volatile *, long); long _InterlockedExchangeAdd_HLERelease(long volatile *, long); __int64 _InterlockedExchangeAdd64_HLEAcquire(__int64 volatile *, __int64); __int64 _InterlockedExchangeAdd64_HLERelease(__int64 volatile *, __int64); -static __inline__ void -__attribute__((__deprecated__("use other intrinsics or C++11 atomics instead"))) -_ReadBarrier(void); -static __inline__ void -__attribute__((__deprecated__("use other intrinsics or C++11 atomics instead"))) -_ReadWriteBarrier(void); +void __attribute__((__deprecated__( + "use other intrinsics or C++11 atomics instead"))) _ReadBarrier(void); +void __attribute__((__deprecated__( + "use other intrinsics or C++11 atomics instead"))) _ReadWriteBarrier(void); unsigned int _rorx_u32(unsigned int, const unsigned int); int _sarx_i32(int, unsigned int); #if __STDC_HOSTED__ @@ -185,9 +163,8 @@ unsigned int _shrx_u32(unsigned int, unsigned int); void _Store_HLERelease(long volatile *, long); void _Store64_HLERelease(__int64 volatile *, __int64); void _StorePointer_HLERelease(void *volatile *, void *); -static __inline__ void -__attribute__((__deprecated__("use other intrinsics or C++11 atomics instead"))) -_WriteBarrier(void); +void __attribute__((__deprecated__( + "use other intrinsics or C++11 atomics instead"))) _WriteBarrier(void); unsigned __int32 xbegin(void); void _xend(void); @@ -197,19 +174,14 @@ void __addgsbyte(unsigned long, unsigned char); void __addgsdword(unsigned long, unsigned long); void __addgsqword(unsigned long, unsigned __int64); void __addgsword(unsigned long, unsigned short); -static __inline__ void __faststorefence(void); void __incgsbyte(unsigned long); void __incgsdword(unsigned long); void __incgsqword(unsigned long); void __incgsword(unsigned long); -static __inline__ void __movsq(unsigned long long *, unsigned long long const *, size_t); -static __inline__ unsigned char __readgsbyte(unsigned long); -static __inline__ unsigned long __readgsdword(unsigned long); -static __inline__ unsigned __int64 __readgsqword(unsigned long); unsigned short __readgsword(unsigned long); unsigned __int64 __shiftleft128(unsigned __int64 _LowPart, @@ -218,7 +190,6 @@ unsigned __int64 __shiftleft128(unsigned __int64 _LowPart, unsigned __int64 __shiftright128(unsigned __int64 _LowPart, unsigned __int64 _HighPart, unsigned char _Shift); -static __inline__ void __stosq(unsigned __int64 *, unsigned __int64, size_t); unsigned char __vmx_on(unsigned __int64 *); unsigned char __vmx_vmclear(unsigned __int64 *); @@ -269,13 +240,9 @@ unsigned __int64 _rorx_u64(unsigned __int64, const unsigned int); __int64 _sarx_i64(__int64, unsigned int); unsigned __int64 _shlx_u64(unsigned __int64, unsigned int); unsigned __int64 _shrx_u64(unsigned __int64, unsigned int); -static __inline__ __int64 __mulh(__int64, __int64); -static __inline__ unsigned __int64 __umulh(unsigned __int64, unsigned __int64); -static __inline__ __int64 _mul128(__int64, __int64, __int64*); -static __inline__ unsigned __int64 _umul128(unsigned __int64, unsigned __int64, unsigned __int64*); @@ -284,29 +251,19 @@ unsigned __int64 _umul128(unsigned __int64, #if defined(__x86_64__) || defined(__arm__) || defined(__aarch64__) -static __inline__ unsigned char _BitScanForward64(unsigned long *_Index, unsigned __int64 _Mask); -static __inline__ unsigned char _BitScanReverse64(unsigned long *_Index, unsigned __int64 _Mask); #endif #if defined(__i386__) || defined(__x86_64__) || defined(__arm__) || defined(__aarch64__) -static __inline__ __int64 _InterlockedDecrement64(__int64 volatile *_Addend); -static __inline__ __int64 _InterlockedExchange64(__int64 volatile *_Target, __int64 _Value); -static __inline__ __int64 _InterlockedExchangeAdd64(__int64 volatile *_Addend, __int64 _Value); -static __inline__ __int64 _InterlockedExchangeSub64(__int64 volatile *_Subend, __int64 _Value); -static __inline__ __int64 _InterlockedIncrement64(__int64 volatile *_Addend); -static __inline__ __int64 _InterlockedOr64(__int64 volatile *_Value, __int64 _Mask); -static __inline__ __int64 _InterlockedXor64(__int64 volatile *_Value, __int64 _Mask); -static __inline__ __int64 _InterlockedAnd64(__int64 volatile *_Value, __int64 _Mask); #endif @@ -475,40 +432,56 @@ __int64 _InterlockedCompareExchange64_rel(__int64 volatile *_Destination, |* movs, stos \*----------------------------------------------------------------------------*/ #if defined(__i386__) || defined(__x86_64__) -static __inline__ void __DEFAULT_FN_ATTRS -__movsb(unsigned char *__dst, unsigned char const *__src, size_t __n) { +static __inline__ void __DEFAULT_FN_ATTRS __movsb(unsigned char *__dst, + unsigned char const *__src, + size_t __n) { __asm__ __volatile__("rep movsb" : "+D"(__dst), "+S"(__src), "+c"(__n) : : "memory"); } -static __inline__ void __DEFAULT_FN_ATTRS -__movsd(unsigned long *__dst, unsigned long const *__src, size_t __n) { - __asm__ __volatile__("rep movsl" : "+D"(__dst), "+S"(__src), "+c"(__n) - : : "memory"); +static __inline__ void __DEFAULT_FN_ATTRS __movsd(unsigned long *__dst, + unsigned long const *__src, + size_t __n) { + __asm__ __volatile__("rep movsl" + : "+D"(__dst), "+S"(__src), "+c"(__n) + : + : "memory"); } -static __inline__ void __DEFAULT_FN_ATTRS -__movsw(unsigned short *__dst, unsigned short const *__src, size_t __n) { - __asm__ __volatile__("rep movsw" : "+D"(__dst), "+S"(__src), "+c"(__n) - : : "memory"); +static __inline__ void __DEFAULT_FN_ATTRS __movsw(unsigned short *__dst, + unsigned short const *__src, + size_t __n) { + __asm__ __volatile__("rep movsw" + : "+D"(__dst), "+S"(__src), "+c"(__n) + : + : "memory"); } -static __inline__ void __DEFAULT_FN_ATTRS -__stosd(unsigned long *__dst, unsigned long __x, size_t __n) { - __asm__ __volatile__("rep stosl" : "+D"(__dst), "+c"(__n) : "a"(__x) +static __inline__ void __DEFAULT_FN_ATTRS __stosd(unsigned long *__dst, + unsigned long __x, + size_t __n) { + __asm__ __volatile__("rep stosl" + : "+D"(__dst), "+c"(__n) + : "a"(__x) : "memory"); } -static __inline__ void __DEFAULT_FN_ATTRS -__stosw(unsigned short *__dst, unsigned short __x, size_t __n) { - __asm__ __volatile__("rep stosw" : "+D"(__dst), "+c"(__n) : "a"(__x) +static __inline__ void __DEFAULT_FN_ATTRS __stosw(unsigned short *__dst, + unsigned short __x, + size_t __n) { + __asm__ __volatile__("rep stosw" + : "+D"(__dst), "+c"(__n) + : "a"(__x) : "memory"); } #endif #ifdef __x86_64__ -static __inline__ void __DEFAULT_FN_ATTRS -__movsq(unsigned long long *__dst, unsigned long long const *__src, size_t __n) { - __asm__ __volatile__("rep movsq" : "+D"(__dst), "+S"(__src), "+c"(__n) - : : "memory"); +static __inline__ void __DEFAULT_FN_ATTRS __movsq( *** 2649 LINES SKIPPED ***
Want to link to this message? Use this URL: <https://mail-archive.FreeBSD.org/cgi/mid.cgi?202101022120.102LK1tB066776>