diff options
Diffstat (limited to 'clang')
23 files changed, 218 insertions, 39 deletions
diff --git a/clang/docs/ClangOffloadBundler.rst b/clang/docs/ClangOffloadBundler.rst index 5570dbb..f63037a7 100644 --- a/clang/docs/ClangOffloadBundler.rst +++ b/clang/docs/ClangOffloadBundler.rst @@ -28,7 +28,7 @@ A bundled code object may also be used to bundle just the offloaded code objects, and embedded as data into the host code object. The host compilation includes an ``init`` function that will use the runtime corresponding to the offload kind (see :ref:`clang-offload-kind-table`) to load the offload code -objects appropriate to the devices present when the host program is executed. +objects appropriate for the devices present when the host program is executed. :program:`clang-offload-bundler` is located in `clang/tools/clang-offload-bundler`. @@ -147,7 +147,7 @@ bundle file is: <end> ::== OFFLOAD_BUNDLER_MAGIC_STR__END__ **comment** - The symbol used for starting single-line comment in the file type of + The symbol used for starting a single-line comment in the file type of constituting bundles. E.g. it is ";" for ll ``File Type`` and "#" for "s" ``File Type``. @@ -155,13 +155,13 @@ bundle file is: The :ref:`clang-bundle-entry-id` for the enclosing bundle. **eol** - The end of line character. + The end-of-line character. **bundle** The code object stored in one of the supported text file formats. **OFFLOAD_BUNDLER_MAGIC_STR__** - Magic string that marks the existence of offloading data i.e. + The magic string that marks the existence of offloading data i.e. "__CLANG_OFFLOAD_BUNDLE__". .. _clang-bundled-code-object-layout: @@ -231,7 +231,7 @@ Where: ============= ============================================================== host Host code object. ``clang-offload-bundler`` always includes this entry as the first bundled code object entry. For an - embedded bundled code object this entry is not used by the + embedded bundled code object, this entry is not used by the runtime and so is generally an empty code object. hip Offload code object for the HIP language. Used for all @@ -272,7 +272,7 @@ without differentiation based on offload kind. ``<arch><sub>-<vendor>-<sys>-<env>`` However, in order to standardize outputs for tools that consume bitcode bundles - and to parse target ID containing dashes, the bundler only accepts target + and to parse a target ID containing dashes, the bundler only accepts target triples in the 4-field format: ``<arch><sub>-<vendor>-<sys>-<env>`` @@ -292,7 +292,7 @@ Bundled Code Object Composition * If there is an entry with a target feature specified as *Any*, then all entries must specify that target feature as *Any* for the same processor. -There may be additional target specific restrictions. +There may be additional target-specific restrictions. .. _compatibility-bundle-entry-id: @@ -300,9 +300,9 @@ Compatibility Rules for Bundle Entry ID --------------------------------------- A code object, specified using its Bundle Entry ID, can be loaded and - executed on a target processor, if: + executed on a target processor if: - * Their offload kinds are the same or comptible. + * Their offload kinds are the same or compatible. * Their target triples are compatible. * Their Target IDs are compatible as defined in :ref:`compatibility-target-id`. @@ -331,7 +331,7 @@ Target ID syntax is defined by the following BNF syntax: Where: **processor** - Is a the target specific processor or any alternative processor name. + Is the target-specific processor or any alternative processor name. **target-feature** Is a target feature name that is supported by the processor. Each target @@ -350,7 +350,7 @@ Where: can only be loaded on a processor configured with the target feature on. *Off* - specified by ``-``, indicating the target feature is disabled. A code + Specified by ``-``, indicating the target feature is disabled. A code object compiled with a target ID specifying a target feature off can only be loaded on a processor configured with the target feature off. @@ -360,9 +360,9 @@ Compatibility Rules for Target ID --------------------------------- A code object compiled for a Target ID is considered compatible for a - target, if: + target if: - * Their processor is same. + * Their processor is the same. * Their feature set is compatible as defined above. There are two forms of target ID: @@ -380,10 +380,10 @@ There are two forms of target ID: alphabetic order. Command line tools convert non-canonical form to canonical form. -Target Specific information +Target-Specific information =========================== -Target specific information is available for the following: +Target-specific information is available for the following: *AMD GPU* AMD GPU supports target ID and target features. See `User Guide for AMDGPU Backend @@ -397,7 +397,7 @@ Most other targets do not support target IDs. Archive Unbundling ================== -Unbundling of a heterogeneous device archive (HDA) is done to create device specific +Unbundling of a heterogeneous device archive (HDA) is done to create device-specific archives. HDA is in a format compatible with GNU ``ar`` utility and contains a collection of bundled device binaries where each bundle file will contain device binaries for a host and one or more targets. The output device-specific @@ -469,7 +469,7 @@ compatible with that particular offload target. Compatibility between a device binary in HDA and a target is based on the compatibility between their bundle entry IDs as defined in :ref:`compatibility-bundle-entry-id`. -Following cases may arise during compatibility testing: +The following cases may arise during compatibility testing: * A binary is compatible with one or more targets: Insert the binary into the device-specific archive of each compatible target. @@ -517,7 +517,7 @@ Compression and Decompression ``clang-offload-bundler`` provides features to compress and decompress the full bundle, leveraging inherent redundancies within the bundle entries. Use the -`-compress` command-line option to enable this compression capability. +``-compress`` command-line option to enable this compression capability. The compressed offload bundle begins with a header followed by the compressed binary data: @@ -542,4 +542,4 @@ The compressed offload bundle begins with a header followed by the compressed bi - **Compressed Data**: The actual compressed binary data follows the header. Its size can be inferred from the total size of the file minus the header size. - > **Note**: Version 3 is now the default format. For backward compatibility with older HIP runtimes that support version 2 only, set the environment variable `COMPRESSED_BUNDLE_FORMAT_VERSION=2`. + > **Note**: Version 3 is now the default format. For backward compatibility with older HIP runtimes that support version 2 only, set the environment variable ``COMPRESSED_BUNDLE_FORMAT_VERSION=2``. diff --git a/clang/docs/SafeBuffers.rst b/clang/docs/SafeBuffers.rst index da75907..205e621 100644 --- a/clang/docs/SafeBuffers.rst +++ b/clang/docs/SafeBuffers.rst @@ -262,7 +262,7 @@ You can achieve this by refactoring the function to accept a ``std::span`` as a parameter:: int get_last_element(std::span<int> sp) { - return sp[size - 1]; + return sp[sp.size() - 1]; } This solution puts the responsibility for making sure the span is well-formed @@ -411,7 +411,7 @@ backwards compatibility -- in terms of both API and ABI -- by adding a "compatibility overload":: int get_last_element(std::span<int> sp) { - return sp[size - 1]; + return sp[sp.size() - 1]; } [[clang::unsafe_buffer_usage]] // Please use the new function. diff --git a/clang/include/clang/ASTMatchers/ASTMatchers.h b/clang/include/clang/ASTMatchers/ASTMatchers.h index 492863d..98e62de 100644 --- a/clang/include/clang/ASTMatchers/ASTMatchers.h +++ b/clang/include/clang/ASTMatchers/ASTMatchers.h @@ -2763,6 +2763,20 @@ extern const internal::VariadicDynCastAllOfMatcher<Stmt, CXXDynamicCastExpr> extern const internal::VariadicDynCastAllOfMatcher<Stmt, CXXConstCastExpr> cxxConstCastExpr; +/// Matches any named cast expression. +/// +/// Example: Matches all four of the casts in +/// \code +/// struct S { virtual void f(); }; +/// S* p = nullptr; +/// S* ptr1 = static_cast<S*>(p); +/// S* ptr2 = reinterpret_cast<S*>(p); +/// S* ptr3 = dynamic_cast<S*>(p); +/// S* ptr4 = const_cast<S*>(p); +/// \endcode +extern const internal::VariadicDynCastAllOfMatcher<Stmt, CXXNamedCastExpr> + cxxNamedCastExpr; + /// Matches a C-style cast expression. /// /// Example: Matches (int) 2.2f in diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 01d121b..f265d82 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -177,15 +177,15 @@ BUILTIN(__builtin_amdgcn_raw_buffer_load_b64, "V2UiQbiiIi", "n") BUILTIN(__builtin_amdgcn_raw_buffer_load_b96, "V3UiQbiiIi", "n") BUILTIN(__builtin_amdgcn_raw_buffer_load_b128, "V4UiQbiiIi", "n") -BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32, "iiQbiiIi", "t") +BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_add_i32, "iiQbiiIi", "") -TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32, "ffQbiiIi", "t", "atomic-fadd-rtn-insts") +TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32, "ffQbiiIi", "", "atomic-fadd-rtn-insts") TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16, "V2hV2hQbiiIi", "t", "atomic-buffer-global-pk-add-f16-insts") -TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32, "ffQbiiIi", "t", "atomic-fmin-fmax-global-f32") -TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32, "ffQbiiIi", "t", "atomic-fmin-fmax-global-f32") -TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64, "ddQbiiIi", "t", "atomic-fmin-fmax-global-f64") -TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64, "ddQbiiIi", "t", "atomic-fmin-fmax-global-f64") +TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32, "ffQbiiIi", "", "atomic-fmin-fmax-global-f32") +TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32, "ffQbiiIi", "", "atomic-fmin-fmax-global-f32") +TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64, "ddQbiiIi", "", "atomic-fmin-fmax-global-f64") +TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64, "ddQbiiIi", "", "atomic-fmin-fmax-global-f64") TARGET_BUILTIN(__builtin_amdgcn_raw_ptr_buffer_load_lds, "vQbv*3IUiiiIiIi", "t", "vmem-to-lds-load-insts") TARGET_BUILTIN(__builtin_amdgcn_struct_ptr_buffer_load_lds, "vQbv*3IUiiiiIiIi", "t", "vmem-to-lds-load-insts") diff --git a/clang/lib/AST/ByteCode/Compiler.cpp b/clang/lib/AST/ByteCode/Compiler.cpp index f4ddbf4..6c08846 100644 --- a/clang/lib/AST/ByteCode/Compiler.cpp +++ b/clang/lib/AST/ByteCode/Compiler.cpp @@ -3948,6 +3948,8 @@ bool Compiler<Emitter>::VisitRecoveryExpr(const RecoveryExpr *E) { template <class Emitter> bool Compiler<Emitter>::VisitAddrLabelExpr(const AddrLabelExpr *E) { assert(E->getType()->isVoidPointerType()); + if (DiscardResult) + return true; return this->emitDummyPtr(E, E); } diff --git a/clang/lib/AST/ByteCode/Interp.cpp b/clang/lib/AST/ByteCode/Interp.cpp index 169a9a2..910868b 100644 --- a/clang/lib/AST/ByteCode/Interp.cpp +++ b/clang/lib/AST/ByteCode/Interp.cpp @@ -832,6 +832,8 @@ bool CheckLoad(InterpState &S, CodePtr OpPC, const Pointer &Ptr, return false; if (!CheckVolatile(S, OpPC, Ptr, AK)) return false; + if (!Ptr.isConst() && !S.inConstantContext() && isConstexprUnknown(Ptr)) + return false; return true; } diff --git a/clang/lib/AST/ByteCode/InterpBuiltin.cpp b/clang/lib/AST/ByteCode/InterpBuiltin.cpp index d0b97a1..839e84f 100644 --- a/clang/lib/AST/ByteCode/InterpBuiltin.cpp +++ b/clang/lib/AST/ByteCode/InterpBuiltin.cpp @@ -1941,6 +1941,9 @@ static bool interp__builtin_memcmp(InterpState &S, CodePtr OpPC, return true; } + if (!PtrA.isBlockPointer() || !PtrB.isBlockPointer()) + return false; + bool IsWide = (ID == Builtin::BIwmemcmp || ID == Builtin::BI__builtin_wmemcmp); diff --git a/clang/lib/ASTMatchers/ASTMatchersInternal.cpp b/clang/lib/ASTMatchers/ASTMatchersInternal.cpp index 1f0e007..42f124b 100644 --- a/clang/lib/ASTMatchers/ASTMatchersInternal.cpp +++ b/clang/lib/ASTMatchers/ASTMatchersInternal.cpp @@ -1009,6 +1009,8 @@ const internal::VariadicDynCastAllOfMatcher<Stmt, CXXDynamicCastExpr> cxxDynamicCastExpr; const internal::VariadicDynCastAllOfMatcher<Stmt, CXXConstCastExpr> cxxConstCastExpr; +const internal::VariadicDynCastAllOfMatcher<Stmt, CXXNamedCastExpr> + cxxNamedCastExpr; const internal::VariadicDynCastAllOfMatcher<Stmt, CStyleCastExpr> cStyleCastExpr; const internal::VariadicDynCastAllOfMatcher<Stmt, ExplicitCastExpr> diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 12e2813ef..6af8066 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -110,6 +110,33 @@ static bool IsArtificial(VarDecl const *VD) { cast<Decl>(VD->getDeclContext())->isImplicit()); } +/// Returns \c true if the specified variable \c VD is an explicit parameter of +/// a synthesized Objective-C property accessor. E.g., a synthesized property +/// setter method will have a single explicit parameter which is the property to +/// set. +static bool IsObjCSynthesizedPropertyExplicitParameter(VarDecl const *VD) { + assert(VD); + + if (!llvm::isa<ParmVarDecl>(VD)) + return false; + + // Not a property method. + const auto *Method = + llvm::dyn_cast_or_null<ObjCMethodDecl>(VD->getDeclContext()); + if (!Method) + return false; + + // Not a synthesized property accessor. + if (!Method->isImplicit() || !Method->isPropertyAccessor()) + return false; + + // Not an explicit parameter. + if (VD->isImplicit()) + return false; + + return true; +} + CGDebugInfo::CGDebugInfo(CodeGenModule &CGM) : CGM(CGM), DebugKind(CGM.getCodeGenOpts().getDebugInfo()), DebugTypeExtRefs(CGM.getCodeGenOpts().DebugTypeExtRefs), @@ -5158,7 +5185,12 @@ llvm::DILocalVariable *CGDebugInfo::EmitDeclare(const VarDecl *VD, } SmallVector<uint64_t, 13> Expr; llvm::DINode::DIFlags Flags = llvm::DINode::FlagZero; - if (VarIsArtificial) + + // While synthesized Objective-C property setters are "artificial" (i.e., they + // are not spelled out in source), we want to pretend they are just like a + // regular non-compiler generated method. Hence, don't mark explicitly passed + // parameters of such methods as artificial. + if (VarIsArtificial && !IsObjCSynthesizedPropertyExplicitParameter(VD)) Flags |= llvm::DINode::FlagArtificial; auto Align = getDeclAlignIfRequired(VD, CGM.getContext()); diff --git a/clang/lib/Interpreter/Interpreter.cpp b/clang/lib/Interpreter/Interpreter.cpp index b05cb5a..cde354c 100644 --- a/clang/lib/Interpreter/Interpreter.cpp +++ b/clang/lib/Interpreter/Interpreter.cpp @@ -581,12 +581,7 @@ Interpreter::Parse(llvm::StringRef Code) { if (!TuOrErr) return TuOrErr.takeError(); - PTUs.emplace_back(PartialTranslationUnit()); - PartialTranslationUnit &LastPTU = PTUs.back(); - LastPTU.TUPart = *TuOrErr; - - if (std::unique_ptr<llvm::Module> M = Act->GenModule()) - LastPTU.TheModule = std::move(M); + PartialTranslationUnit &LastPTU = IncrParser->RegisterPTU(*TuOrErr); return LastPTU; } diff --git a/clang/test/AST/ByteCode/builtin-functions.cpp b/clang/test/AST/ByteCode/builtin-functions.cpp index 0b7d51b..d8572ba 100644 --- a/clang/test/AST/ByteCode/builtin-functions.cpp +++ b/clang/test/AST/ByteCode/builtin-functions.cpp @@ -1510,6 +1510,8 @@ namespace Memcmp { static_assert(f()); #endif + int unknown; + void foo(void) { unknown *= __builtin_memcmp(0, 0, 2); } } namespace Memchr { diff --git a/clang/test/AST/ByteCode/codegen-cxx20.cpp b/clang/test/AST/ByteCode/codegen-cxx20.cpp new file mode 100644 index 0000000..c1ef629 --- /dev/null +++ b/clang/test/AST/ByteCode/codegen-cxx20.cpp @@ -0,0 +1,15 @@ +// RUN: %clang_cc1 -triple x86_64-linux -emit-llvm -o - %s -fcxx-exceptions | FileCheck %s +// RUN: %clang_cc1 -triple x86_64-linux -emit-llvm -o - %s -fcxx-exceptions -fexperimental-new-constant-interpreter | FileCheck %s + + +/// The read from a used to succeed, causing the entire if statement to vanish. +extern void e(); +int somefunc() { + auto foo = [a = false]() mutable { + if (a) + e(); + }; + foo(); +} + +// CHECK: call void @_Z1ev() diff --git a/clang/test/AST/ByteCode/cxx11.cpp b/clang/test/AST/ByteCode/cxx11.cpp index 8efd320..427d3a1 100644 --- a/clang/test/AST/ByteCode/cxx11.cpp +++ b/clang/test/AST/ByteCode/cxx11.cpp @@ -370,3 +370,12 @@ namespace GH150709 { static_assert((e2[0].*mp)() == 1, ""); // ref-error {{constant expression}} static_assert((g.*mp)() == 1, ""); // ref-error {{constant expression}} } + +namespace DiscardedAddrLabel { + void foo(void) { + L: + *&&L; // both-error {{indirection not permitted}} \ + // both-warning {{expression result unused}} + } +} + diff --git a/clang/test/ClangScanDeps/resource_directory.c b/clang/test/ClangScanDeps/resource_directory.c index 6183e8a..5c4b24f 100644 --- a/clang/test/ClangScanDeps/resource_directory.c +++ b/clang/test/ClangScanDeps/resource_directory.c @@ -1,4 +1,5 @@ -// REQUIRES: shell +// Path seperator differences +// UNSUPPORTED: system-windows // RUN: rm -rf %t && mkdir %t // RUN: cp %S/Inputs/resource_directory/* %t diff --git a/clang/test/CodeGen/arm-target-features.c b/clang/test/CodeGen/arm-target-features.c index 95ae27bd..2b5a410 100644 --- a/clang/test/CodeGen/arm-target-features.c +++ b/clang/test/CodeGen/arm-target-features.c @@ -116,6 +116,9 @@ // RUN: %clang_cc1 -triple thumb-linux-gnueabi -target-cpu cortex-m52 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-ARMV81M-CORTEX-M52-LINUX // CHECK-ARMV81M-CORTEX-M52-LINUX: "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+pacbti,+ras,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp" +// RUN: %clang_cc1 -triple thumb-linux-gnueabi -target-cpu star-mc3 -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-ARMV81M-STAR-MC3-LINUX +// CHECK-ARMV81M-STAR-MC3-LINUX: "target-features"="+armv8.1-m.main,+dsp,+fp-armv8d16,+fp-armv8d16sp,+fp16,+fp64,+fullfp16,+hwdiv,+lob,+mve,+mve.fp,+pacbti,+ras,+thumb-mode,+vfp2,+vfp2sp,+vfp3d16,+vfp3d16sp,+vfp4d16,+vfp4d16sp" + // RUN: %clang_cc1 -triple thumbv9.3a-linux-gnueabihf -emit-llvm -o - %s | FileCheck %s --check-prefix=CHECK-ARCH93 // CHECK-ARCH93: "target-features"="+armv9.3-a,+thumb-mode,+v8.1a,+v8.2a,+v8.3a,+v8.4a,+v8.5a,+v8.6a,+v8.7a,+v8.8a,+v9.1a,+v9.2a,+v9.3a,+v9a" diff --git a/clang/test/DebugInfo/ObjC/property-synthesized-accessors.m b/clang/test/DebugInfo/ObjC/property-synthesized-accessors.m new file mode 100644 index 0000000..d2e2dba --- /dev/null +++ b/clang/test/DebugInfo/ObjC/property-synthesized-accessors.m @@ -0,0 +1,63 @@ +// Test that synthesized accessors get treated like regular method declarations/definitions. +// I.e.: +// 1. explicitly passed parameter are not marked artificial. +// 2. Each property accessor has a method declaration and definition. + +// RUN: %clang_cc1 -triple %itanium_abi_triple -emit-llvm -dwarf-version=5 -debug-info-kind=limited %s -o - | FileCheck %s --implicit-check-not "DIFlagArtificial" + +@interface Foo +@property int p1; +@end + +@implementation Foo +@end + +int main(void) { + Foo *f; + f.p1 = 2; + return f.p1; +} + +// CHECK: ![[P1_TYPE:[0-9]+]] = !DIBasicType(name: "int" +// CHECK: ![[GETTER_DECL:[0-9]+]] = !DISubprogram(name: "-[Foo p1]" +// CHECK-SAME: type: ![[GETTER_TYPE:[0-9]+]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped +// CHECK-SAME: spFlags: DISPFlagLocalToUnit) + +// CHECK: ![[GETTER_TYPE]] = !DISubroutineType(types: ![[GETTER_PARAMS:[0-9]+]]) +// CHECK: ![[GETTER_PARAMS]] = !{![[P1_TYPE]], ![[ID_TYPE:[0-9]+]], ![[SEL_TYPE:[0-9]+]]} +// CHECK: ![[ID_TYPE]] = !DIDerivedType(tag: DW_TAG_pointer_type +// CHECK-SAME: flags: DIFlagArtificial | DIFlagObjectPointer) +// CHECK: ![[SEL_TYPE]] = !DIDerivedType(tag: DW_TAG_typedef, name: "SEL" +// CHECK-SAME: flags: DIFlagArtificial) + +// CHECK: ![[SETTER_DECL:[0-9]+]] = !DISubprogram(name: "-[Foo setP1:]" +// CHECK-SAME: type: ![[SETTER_TYPE:[0-9]+]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped +// CHECK-SAME: spFlags: DISPFlagLocalToUnit) +// CHECK: ![[SETTER_TYPE]] = !DISubroutineType(types: ![[SETTER_PARAMS:[0-9]+]]) +// CHECK: ![[SETTER_PARAMS]] = !{null, ![[ID_TYPE]], ![[SEL_TYPE]], ![[P1_TYPE]]} + +// CHECK: ![[GETTER_DEF:[0-9]+]] = distinct !DISubprogram(name: "-[Foo p1]" +// CHECK-SAME: type: ![[GETTER_TYPE]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped +// CHECK-SAME: spFlags: DISPFlagLocalToUnit | DISPFlagDefinition +// CHECK-SAME: declaration: ![[GETTER_DECL]] + +// CHECK: !DILocalVariable(name: "self", arg: 1, scope: ![[GETTER_DEF]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagObjectPointer) +// +// CHECK: !DILocalVariable(name: "_cmd", arg: 2, scope: ![[GETTER_DEF]], +// CHECK-SAME: flags: DIFlagArtificial) + +// CHECK: ![[SETTER_DEF:[0-9]+]] = distinct !DISubprogram(name: "-[Foo setP1:]", +// CHECK-SAME: type: ![[SETTER_TYPE]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagPrototyped +// CHECK-SAME: spFlags: DISPFlagLocalToUnit | DISPFlagDefinition +// CHECK-SAME: declaration: ![[SETTER_DECL]] + +// CHECK: !DILocalVariable(name: "self", arg: 1, scope: ![[SETTER_DEF]] +// CHECK-SAME: flags: DIFlagArtificial | DIFlagObjectPointer +// CHECK: !DILocalVariable(name: "_cmd", arg: 2, scope: ![[SETTER_DEF]] +// CHECK-SAME: flags: DIFlagArtificial +// CHECK: !DILocalVariable(name: "p1", arg: 3, scope: ![[SETTER_DEF]] diff --git a/clang/test/Driver/arm-cortex-cpus-2.c b/clang/test/Driver/arm-cortex-cpus-2.c index 0ee8e05..ecdf530 100644 --- a/clang/test/Driver/arm-cortex-cpus-2.c +++ b/clang/test/Driver/arm-cortex-cpus-2.c @@ -585,6 +585,9 @@ // RUN: %clang -target arm -mcpu=cortex-m52 -### -c %s 2>&1 | FileCheck -check-prefix=CHECK-CORTEX-M52 %s // CHECK-CORTEX-M52: "-cc1"{{.*}} "-triple" "thumbv8.1m.main-{{.*}} "-target-cpu" "cortex-m52" +// RUN: %clang -target arm -mcpu=star-mc3 -### -c %s 2>&1 | FileCheck -check-prefix=CHECK-STAR-MC3 %s +// CHECK-STAR-MC3: "-cc1"{{.*}} "-triple" "thumbv8.1m.main-{{.*}} "-target-cpu" "star-mc3" + // RUN: %clang -target arm -mcpu=neoverse-n2 -### -c %s 2>&1 | FileCheck -check-prefix=CHECK-NEOVERSE-N2 %s // CHECK-NEOVERSE-N2: "-cc1"{{.*}} "-triple" "armv9a-{{.*}}" "-target-cpu" "neoverse-n2" diff --git a/clang/test/Driver/baremetal-multilib-custom-error.yaml b/clang/test/Driver/baremetal-multilib-custom-error.yaml index 0be92e2..bc06ed4 100644 --- a/clang/test/Driver/baremetal-multilib-custom-error.yaml +++ b/clang/test/Driver/baremetal-multilib-custom-error.yaml @@ -1,4 +1,3 @@ -# REQUIRES: shell # UNSUPPORTED: system-windows # RUN: %clang --multi-lib-config=%s -no-canonical-prefixes -print-multi-directory 2>&1 \ diff --git a/clang/test/Frontend/absolute-paths-symlinks.c b/clang/test/Frontend/absolute-paths-symlinks.c index 8170910..80bca34 100644 --- a/clang/test/Frontend/absolute-paths-symlinks.c +++ b/clang/test/Frontend/absolute-paths-symlinks.c @@ -12,6 +12,5 @@ // CHECK-SAME: error: unknown type name This do not compile -// REQUIRES: shell // Don't make symlinks on Windows. // UNSUPPORTED: system-windows diff --git a/clang/test/Misc/target-invalid-cpu-note/arm.c b/clang/test/Misc/target-invalid-cpu-note/arm.c index 12acdab..8ac0ed7 100644 --- a/clang/test/Misc/target-invalid-cpu-note/arm.c +++ b/clang/test/Misc/target-invalid-cpu-note/arm.c @@ -70,6 +70,7 @@ // CHECK-SAME: {{^}}, cortex-m55 // CHECK-SAME: {{^}}, cortex-m85 // CHECK-SAME: {{^}}, cortex-m52 +// CHECK-SAME: {{^}}, star-mc3 // CHECK-SAME: {{^}}, cortex-a32 // CHECK-SAME: {{^}}, cortex-a35 // CHECK-SAME: {{^}}, cortex-a53 diff --git a/clang/test/SemaHIP/builtins-amdgcn-raw-buffer-atomic-add.hip b/clang/test/SemaHIP/builtins-amdgcn-raw-buffer-atomic-add.hip new file mode 100644 index 0000000..8ee64d4 --- /dev/null +++ b/clang/test/SemaHIP/builtins-amdgcn-raw-buffer-atomic-add.hip @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx90a -verify %s -fcuda-is-device +// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify %s + +typedef _Float16 __attribute__((ext_vector_type(2))) float16x2_t; + +#define __device__ __attribute__((device)) + +__device__ void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, int i32, float f32, float16x2_t v2f16, int offset, int soffset) { + i32 = __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(i32, rsrc, offset, soffset, 0); + f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(f32, rsrc, offset, soffset, 0); + v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(v2f16, rsrc, offset, soffset, 0); +} + +__device__ void test_raw_ptr_atomics_err(__amdgpu_buffer_rsrc_t rsrc, int i32, float f32, float16x2_t v2f16, int offset, int soffset) { + i32 = __builtin_amdgcn_raw_ptr_buffer_atomic_add_i32(i32, rsrc, offset, soffset, 0, 4); // expected-error{{too many arguments to function call}} + f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_f32(f32, rsrc, offset, soffset, 0, 4); // expected-error{{too many arguments to function call}} + v2f16 = __builtin_amdgcn_raw_ptr_buffer_atomic_fadd_v2f16(v2f16, rsrc, offset, soffset, 0, 4); +} diff --git a/clang/test/SemaHIP/builtins-amdgcn-raw-buffer-atomic-fmin-max.hip b/clang/test/SemaHIP/builtins-amdgcn-raw-buffer-atomic-fmin-max.hip new file mode 100644 index 0000000..a2dc021 --- /dev/null +++ b/clang/test/SemaHIP/builtins-amdgcn-raw-buffer-atomic-fmin-max.hip @@ -0,0 +1,18 @@ +// RUN: %clang_cc1 -fsyntax-only -triple amdgcn -target-cpu gfx90a -verify %s -fcuda-is-device +// RUN: %clang_cc1 -fsyntax-only -triple x86_64 -aux-triple amdgcn -verify %s + +#define __device__ __attribute__((device)) + +__device__ void test_raw_ptr_atomics(__amdgpu_buffer_rsrc_t rsrc, float f32, double f64, int offset, int soffset) { + f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32(f32, rsrc, offset, soffset, 0); + f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64(f64, rsrc, offset, soffset, 0); + f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset, soffset, 0); + f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset, soffset, 0); +} + +__device__ void test_raw_ptr_atomics_err(__amdgpu_buffer_rsrc_t rsrc, float f32, double f64, int offset, int soffset) { + f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f32(f32, rsrc, offset, soffset, 0, 4); // expected-error{{too many arguments to function call}} + f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmin_f64(f64, rsrc, offset, soffset, 0, 4); // expected-error{{too many arguments to function call}} + f32 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f32(f32, rsrc, offset, soffset, 0, 4); // expected-error{{too many arguments to function call}} + f64 = __builtin_amdgcn_raw_ptr_buffer_atomic_fmax_f64(f64, rsrc, offset, soffset, 0, 4); // expected-error{{too many arguments to function call}} +} diff --git a/clang/test/Tooling/clang-check-pwd.cpp b/clang/test/Tooling/clang-check-pwd.cpp index 309cee5..e4360c0 100644 --- a/clang/test/Tooling/clang-check-pwd.cpp +++ b/clang/test/Tooling/clang-check-pwd.cpp @@ -12,5 +12,3 @@ // CHECK: a type specifier is required // CHECK: .foobar/test.cpp invalid; - -// REQUIRES: shell |