diff options
Diffstat (limited to 'clang/include')
29 files changed, 729 insertions, 108 deletions
diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 0273109..3b98274a 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -1818,7 +1818,7 @@ public: NumPositiveBits = std::max({NumPositiveBits, ActiveBits, 1u}); } else { NumNegativeBits = - std::max(NumNegativeBits, (unsigned)InitVal.getSignificantBits()); + std::max(NumNegativeBits, InitVal.getSignificantBits()); } MembersRepresentableByInt &= isRepresentableIntegerValue(InitVal, IntTy); diff --git a/clang/include/clang/AST/Expr.h b/clang/include/clang/AST/Expr.h index 708c6e2..6c124aa 100644 --- a/clang/include/clang/AST/Expr.h +++ b/clang/include/clang/AST/Expr.h @@ -16,7 +16,6 @@ #include "clang/AST/APNumericStorage.h" #include "clang/AST/APValue.h" #include "clang/AST/ASTVector.h" -#include "clang/AST/Attr.h" #include "clang/AST/ComputeDependence.h" #include "clang/AST/Decl.h" #include "clang/AST/DeclAccessPair.h" @@ -58,6 +57,7 @@ namespace clang { class StringLiteral; class TargetInfo; class ValueDecl; + class WarnUnusedResultAttr; /// A simple array of base specifiers. typedef SmallVector<CXXBaseSpecifier*, 4> CXXCastPath; diff --git a/clang/include/clang/Basic/Attr.td b/clang/include/clang/Basic/Attr.td index 224cb6a..b3ff45b 100644 --- a/clang/include/clang/Basic/Attr.td +++ b/clang/include/clang/Basic/Attr.td @@ -1642,7 +1642,7 @@ def DeviceKernel : DeclOrTypeAttr { } def SYCLKernelEntryPoint : InheritableAttr { - let Spellings = [Clang<"sycl_kernel_entry_point">]; + let Spellings = [CXX11<"clang", "sycl_kernel_entry_point">]; let Args = [ // KernelName is required and specifies the kernel name type. TypeArgument<"KernelName">, diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 945e11b..bb3953e 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -651,6 +651,16 @@ TARGET_BUILTIN(__builtin_amdgcn_global_load_monitor_b128, "V4iV4i*1Ii", "nc", "g TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b32, "ii*0Ii", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b64, "V2iV2i*0Ii", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_flat_load_monitor_b128, "V4iV4i*0Ii", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b8, "vc*1c*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b32, "vi*1i*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b64, "vV2i*1V2i*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_load_async_to_lds_b128, "vV4i*1V4i*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b8, "vc*1c*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b32, "vi*1i*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b64, "vV2i*1V2i*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_global_store_async_from_lds_b128, "vV4i*1V4i*3IiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64, "vLi*3", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64, "LiLi*3Li", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_tensor_load_to_lds, "vV4iV8iV4iV4iIi", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_tensor_load_to_lds_d2, "vV4iV8iIi", "nc", "gfx1250-insts") @@ -670,9 +680,6 @@ TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8i16, "V8sV8s*3", "nc", "gfx1 TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8f16, "V8hV8h*3", "nc", "gfx1250-insts,wavefrontsize32") TARGET_BUILTIN(__builtin_amdgcn_ds_load_tr16_b128_v8bf16, "V8yV8y*3", "nc", "gfx1250-insts,wavefrontsize32") -TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_async_barrier_arrive_b64, "vLi*3", "nc", "gfx1250-insts") -TARGET_BUILTIN(__builtin_amdgcn_ds_atomic_barrier_arrive_rtn_b64, "LiLi*3Li", "nc", "gfx1250-insts") - TARGET_BUILTIN(__builtin_amdgcn_s_setprio_inc_wg, "vIs", "n", "setprio-inc-wg-inst") TARGET_BUILTIN(__builtin_amdgcn_s_monitor_sleep, "vIs", "n", "gfx1250-insts") @@ -690,10 +697,27 @@ TARGET_BUILTIN(__builtin_amdgcn_exp2_bf16, "yy", "nc", "bf16-trans-insts") TARGET_BUILTIN(__builtin_amdgcn_sin_bf16, "yy", "nc", "bf16-trans-insts") TARGET_BUILTIN(__builtin_amdgcn_cos_bf16, "yy", "nc", "bf16-trans-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_pk_bf16_f32, "V2yffi", "nc", "bf16-cvt-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_pk_f16_f32, "V2hffi", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_fp8, "hiIi", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_cvt_f16_bf8, "hiIi", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_fp8, "V2hs", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_f16_bf8, "V2hs", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f16, "sV2h", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_bf8_f16, "sV2h", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f16, "ihiUiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_bf8_f16, "ihiUiIi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_f16_fp8, "V8hV2UiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_bf16_fp8, "V8yV2UiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_f16_bf8, "V8hV2UiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_bf16_bf8, "V8yV2UiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_f16_fp4, "V8hUiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_bf16_fp4, "V8yUiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_f32_fp8, "V8fV2UiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_f32_bf8, "V8fV2UiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_scale_pk8_f32_fp4, "V8fUiUiIUi", "nc", "gfx1250-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_pk_fp8_f32_e5m3, "iffiIb", "nc", "fp8e5m3-insts") +TARGET_BUILTIN(__builtin_amdgcn_cvt_sr_fp8_f32_e5m3, "ifiiIi", "nc", "fp8e5m3-insts") TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_i4_i8, "UsUi", "nc", "gfx1250-insts") TARGET_BUILTIN(__builtin_amdgcn_sat_pk4_u4_u8, "UsUi", "nc", "gfx1250-insts") diff --git a/clang/include/clang/Basic/BuiltinsNVPTX.td b/clang/include/clang/Basic/BuiltinsNVPTX.td index 6e531ef..2d6fa17 100644 --- a/clang/include/clang/Basic/BuiltinsNVPTX.td +++ b/clang/include/clang/Basic/BuiltinsNVPTX.td @@ -21,13 +21,17 @@ class SM<string version, list<SMFeatures> newer_list> : SMFeatures { !strconcat(f, "|", newer.Features)); } +let Features = "sm_121a" in def SM_121a : SMFeatures; let Features = "sm_120a" in def SM_120a : SMFeatures; +let Features = "sm_103a" in def SM_103a : SMFeatures; let Features = "sm_101a" in def SM_101a : SMFeatures; let Features = "sm_100a" in def SM_100a : SMFeatures; let Features = "sm_90a" in def SM_90a : SMFeatures; -def SM_120 : SM<"120", [SM_120a]>; -def SM_101 : SM<"101", [SM_101a, SM_120]>; +def SM_121 : SM<"121", [SM_121a]>; +def SM_120 : SM<"120", [SM_120a, SM_121]>; +def SM_103 : SM<"103", [SM_103a, SM_120]>; +def SM_101 : SM<"101", [SM_101a, SM_103]>; def SM_100 : SM<"100", [SM_100a, SM_101]>; def SM_90 : SM<"90", [SM_90a, SM_100]>; def SM_89 : SM<"89", [SM_90]>; @@ -50,8 +54,9 @@ class PTX<string version, PTXFeatures newer> : PTXFeatures { let Features = !strconcat("ptx", version, "|", newer.Features); } -let Features = "ptx87" in def PTX87 : PTXFeatures; +let Features = "ptx88" in def PTX88 : PTXFeatures; +def PTX87 : PTX<"87", PTX88>; def PTX86 : PTX<"86", PTX87>; def PTX85 : PTX<"85", PTX86>; def PTX84 : PTX<"84", PTX85>; diff --git a/clang/include/clang/Basic/Cuda.h b/clang/include/clang/Basic/Cuda.h index d6a22a7..81a792d 100644 --- a/clang/include/clang/Basic/Cuda.h +++ b/clang/include/clang/Basic/Cuda.h @@ -47,9 +47,10 @@ enum class CudaVersion { CUDA_125, CUDA_126, CUDA_128, - FULLY_SUPPORTED = CUDA_123, + CUDA_129, + FULLY_SUPPORTED = CUDA_128, PARTIALLY_SUPPORTED = - CUDA_128, // Partially supported. Proceed with a warning. + CUDA_129, // Partially supported. Proceed with a warning. NEW = 10000, // Too new. Issue a warning, but allow using it. }; const char *CudaVersionToString(CudaVersion V); diff --git a/clang/include/clang/Basic/CustomizableOptional.h b/clang/include/clang/Basic/CustomizableOptional.h index 2d6ae6a..8559eaa 100644 --- a/clang/include/clang/Basic/CustomizableOptional.h +++ b/clang/include/clang/Basic/CustomizableOptional.h @@ -70,15 +70,6 @@ public: void reset() { Storage.reset(); } - LLVM_DEPRECATED("Use &*X instead.", "&*X") - constexpr const T *getPointer() const { return &Storage.value(); } - LLVM_DEPRECATED("Use &*X instead.", "&*X") - T *getPointer() { return &Storage.value(); } - LLVM_DEPRECATED("std::optional::value is throwing. Use *X instead", "*X") - constexpr const T &value() const & { return Storage.value(); } - LLVM_DEPRECATED("std::optional::value is throwing. Use *X instead", "*X") - T &value() & { return Storage.value(); } - constexpr explicit operator bool() const { return has_value(); } constexpr bool has_value() const { return Storage.has_value(); } constexpr const T *operator->() const { return &Storage.value(); } @@ -90,8 +81,6 @@ public: return has_value() ? operator*() : std::forward<U>(alt); } - LLVM_DEPRECATED("std::optional::value is throwing. Use *X instead", "*X") - T &&value() && { return std::move(Storage.value()); } T &&operator*() && { return std::move(Storage.value()); } template <typename U> T value_or(U &&alt) && { diff --git a/clang/include/clang/Basic/DiagnosticASTKinds.td b/clang/include/clang/Basic/DiagnosticASTKinds.td index 071a38f..46d04b0 100644 --- a/clang/include/clang/Basic/DiagnosticASTKinds.td +++ b/clang/include/clang/Basic/DiagnosticASTKinds.td @@ -511,6 +511,14 @@ def note_odr_number_of_bases : Note< "class has %0 base %plural{1:class|:classes}0">; def note_odr_enumerator : Note<"enumerator %0 with value %1 here">; def note_odr_missing_enumerator : Note<"no corresponding enumerator here">; +def note_odr_incompatible_fixed_underlying_type : Note< + "enumeration %0 declared with incompatible fixed underlying types (%1 vs. " + "%2)">; +def note_odr_fixed_underlying_type : Note< + "enumeration %0 has fixed underlying type here">; +def note_odr_missing_fixed_underlying_type : Note< + "enumeration %0 missing fixed underlying type here">; + def err_odr_field_type_inconsistent : Error< "field %0 declared with incompatible types in different " "translation units (%1 vs. %2)">; diff --git a/clang/include/clang/Basic/DiagnosticIDs.h b/clang/include/clang/Basic/DiagnosticIDs.h index f07a003..b21a3b6 100644 --- a/clang/include/clang/Basic/DiagnosticIDs.h +++ b/clang/include/clang/Basic/DiagnosticIDs.h @@ -272,6 +272,11 @@ public: DiagnosticIDs(); ~DiagnosticIDs(); + // Convenience method to construct a new refcounted DiagnosticIDs. + static llvm::IntrusiveRefCntPtr<DiagnosticIDs> create() { + return llvm::makeIntrusiveRefCnt<DiagnosticIDs>(); + } + /// Return an ID for a diagnostic with the specified format string and /// level. /// diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index da4216d..94b174c 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -1812,10 +1812,7 @@ def note_unsatisfied_trait_reason "%DeletedAssign{has a deleted %select{copy|move}1 " "assignment operator}|" "%UnionWithUserDeclaredSMF{is a union with a user-declared " - "%sub{select_special_member_kind}1}|" - "%FunctionType{is a function type}|" - "%CVVoidType{is a cv void type}|" - "%IncompleteArrayType{is an incomplete array type}" + "%sub{select_special_member_kind}1}" "}0">; def warn_consteval_if_always_true : Warning< @@ -9331,8 +9328,28 @@ def err_atomic_builtin_pointer_size : Error< "address argument to atomic builtin must be a pointer to 1,2,4,8 or 16 byte " "type (%0 invalid)">; def err_atomic_exclusive_builtin_pointer_size : Error< - "address argument to load or store exclusive builtin must be a pointer to" - " 1,2,4 or 8 byte type (%0 invalid)">; + "address argument to load or store exclusive builtin must be a pointer to " + // Because the range of legal sizes for load/store exclusive varies with the + // Arm architecture version, this error message wants to be able to specify + // various different subsets of the sizes 1, 2, 4, 8. Rather than make a + // separate diagnostic for each subset, I've arranged here that _this_ error + // can display any combination of the sizes. For each size there are two + // %select parameters: the first chooses whether you need a "," or " or " to + // separate the number from a previous one (or neither), and the second + // parameter indicates whether to display the number itself. + // + // (The very first of these parameters isn't really necessary, since you + // never want to start with "," or " or " before the first number in the + // list, but it keeps it simple to make it look exactly like the other cases, + // and also allows a loop constructing this diagnostic to handle every case + // exactly the same.) + "%select{|,| or }1%select{|1}2" + "%select{|,| or }3%select{|2}4" + "%select{|,| or }5%select{|4}6" + "%select{|,| or }7%select{|8}8" + " byte type (%0 invalid)">; +def err_atomic_exclusive_builtin_pointer_size_none : Error< + "load and store exclusive builtins are not available on this architecture">; def err_atomic_builtin_ext_int_size : Error< "atomic memory operand must have a power-of-two size">; def err_atomic_builtin_bit_int_prohibit : Error< @@ -12919,31 +12936,29 @@ def err_sycl_special_type_num_init_method : Error< // SYCL kernel entry point diagnostics def err_sycl_entry_point_invalid : Error< - "'sycl_kernel_entry_point' attribute cannot be applied to a" + "the %0 attribute cannot be applied to a" " %select{non-static member function|variadic function|deleted function|" "defaulted function|constexpr function|consteval function|" "function declared with the 'noreturn' attribute|coroutine|" - "function defined with a function try block}0">; + "function defined with a function try block}1">; def err_sycl_entry_point_invalid_redeclaration : Error< - "'sycl_kernel_entry_point' kernel name argument does not match prior" - " declaration%diff{: $ vs $|}0,1">; + "the %0 kernel name argument does not match prior" + " declaration%diff{: $ vs $|}1,2">; def err_sycl_kernel_name_conflict : Error< - "'sycl_kernel_entry_point' kernel name argument conflicts with a previous" - " declaration">; + "the %0 kernel name argument conflicts with a previous declaration">; def warn_sycl_kernel_name_not_a_class_type : Warning< "%0 is not a valid SYCL kernel name type; a non-union class type is required">, InGroup<DiagGroup<"nonportable-sycl">>, DefaultError; def warn_sycl_entry_point_redundant_declaration : Warning< - "redundant 'sycl_kernel_entry_point' attribute">, InGroup<RedundantAttribute>; + "redundant %0 attribute">, InGroup<RedundantAttribute>; def err_sycl_entry_point_after_definition : Error< - "'sycl_kernel_entry_point' attribute cannot be added to a function after the" - " function is defined">; + "the %0 attribute cannot be added to a function after the function is" + " defined">; def err_sycl_entry_point_return_type : Error< - "'sycl_kernel_entry_point' attribute only applies to functions with a" - " 'void' return type">; + "the %0 attribute only applies to functions with a 'void' return type">; def err_sycl_entry_point_deduced_return_type : Error< - "'sycl_kernel_entry_point' attribute only applies to functions with a" - " non-deduced 'void' return type">; + "the %0 attribute only applies to functions with a non-deduced 'void' return" + " type">; def warn_cuda_maxclusterrank_sm_90 : Warning< "maxclusterrank requires sm_90 or higher, CUDA arch provided: %0, ignoring " diff --git a/clang/include/clang/Basic/FileManager.h b/clang/include/clang/Basic/FileManager.h index e83a61d..337911e 100644 --- a/clang/include/clang/Basic/FileManager.h +++ b/clang/include/clang/Basic/FileManager.h @@ -237,11 +237,6 @@ public: FileEntryRef getVirtualFileRef(StringRef Filename, off_t Size, time_t ModificationTime); - LLVM_DEPRECATED("Functions returning FileEntry are deprecated.", - "getVirtualFileRef()") - const FileEntry *getVirtualFile(StringRef Filename, off_t Size, - time_t ModificationTime); - /// Retrieve a FileEntry that bypasses VFE, which is expected to be a virtual /// file entry, to access the real file. The returned FileEntry will have /// the same filename as FE but a different identity and its own stat. diff --git a/clang/include/clang/Basic/OffloadArch.h b/clang/include/clang/Basic/OffloadArch.h index 4dda3ec..387a684 100644 --- a/clang/include/clang/Basic/OffloadArch.h +++ b/clang/include/clang/Basic/OffloadArch.h @@ -45,8 +45,12 @@ enum class OffloadArch { SM_100a, SM_101, SM_101a, + SM_103, + SM_103a, SM_120, SM_120a, + SM_121, + SM_121a, GFX600, GFX601, GFX602, diff --git a/clang/include/clang/Basic/TargetInfo.h b/clang/include/clang/Basic/TargetInfo.h index abfbdfa..ce4677e 100644 --- a/clang/include/clang/Basic/TargetInfo.h +++ b/clang/include/clang/Basic/TargetInfo.h @@ -1071,6 +1071,17 @@ public: /// as Custom Datapath. uint32_t getARMCDECoprocMask() const { return ARMCDECoprocMask; } + /// For ARM targets returns a mask defining which data sizes are suitable for + /// __builtin_arm_ldrex and __builtin_arm_strex. + enum { + ARM_LDREX_B = (1 << 0), /// byte (8-bit) + ARM_LDREX_H = (1 << 1), /// half (16-bit) + ARM_LDREX_W = (1 << 2), /// word (32-bit) + ARM_LDREX_D = (1 << 3), /// double (64-bit) + }; + + virtual unsigned getARMLDREXMask() const { return 0; } + /// Returns whether the passed in string is a valid clobber in an /// inline asm statement. /// diff --git a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h index 5c04d59..b26e558 100644 --- a/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h +++ b/clang/include/clang/CIR/Dialect/Builder/CIRBaseBuilder.h @@ -75,6 +75,12 @@ public: return getConstant(loc, cir::IntAttr::get(ty, value)); } + mlir::Value getSignedInt(mlir::Location loc, int64_t val, unsigned numBits) { + auto type = cir::IntType::get(getContext(), numBits, /*isSigned=*/true); + return getConstAPInt(loc, type, + llvm::APInt(numBits, val, /*isSigned=*/true)); + } + mlir::Value getUnsignedInt(mlir::Location loc, uint64_t val, unsigned numBits) { auto type = cir::IntType::get(getContext(), numBits, /*isSigned=*/false); @@ -441,6 +447,10 @@ public: return create<cir::CmpOp>(loc, getBoolTy(), kind, lhs, rhs); } + mlir::Value createIsNaN(mlir::Location loc, mlir::Value operand) { + return createCompare(loc, cir::CmpOpKind::ne, operand, operand); + } + mlir::Value createShift(mlir::Location loc, mlir::Value lhs, mlir::Value rhs, bool isShiftLeft) { return create<cir::ShiftOp>(loc, lhs.getType(), lhs, rhs, isShiftLeft); diff --git a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td index 9961544..588fb0d 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td +++ b/clang/include/clang/CIR/Dialect/IR/CIRAttrs.td @@ -42,7 +42,7 @@ class CIR_TypedAttr<string name, string attrMnemonic, list<Trait> traits = []> let assemblyFormat = [{}]; } -class CIRUnitAttr<string name, string attrMnemonic, list<Trait> traits = []> +class CIR_UnitAttr<string name, string attrMnemonic, list<Trait> traits = []> : CIR_Attr<name, attrMnemonic, traits> { let returnType = "bool"; let defaultValue = "false"; @@ -127,7 +127,7 @@ def CIR_BoolAttr : CIR_Attr<"Bool", "bool", [TypedAttrInterface]> { // ZeroAttr //===----------------------------------------------------------------------===// -def ZeroAttr : CIR_TypedAttr<"Zero", "zero"> { +def CIR_ZeroAttr : CIR_TypedAttr<"Zero", "zero"> { let summary = "Attribute to represent zero initialization"; let description = [{ The ZeroAttr is used to indicate zero initialization on structs. @@ -138,7 +138,7 @@ def ZeroAttr : CIR_TypedAttr<"Zero", "zero"> { // UndefAttr //===----------------------------------------------------------------------===// -def UndefAttr : CIR_TypedAttr<"Undef", "undef"> { +def CIR_UndefAttr : CIR_TypedAttr<"Undef", "undef"> { let summary = "Represent an undef constant"; let description = [{ The UndefAttr represents an undef constant, corresponding to LLVM's notion @@ -264,7 +264,9 @@ def CIR_FPAttr : CIR_Attr<"FP", "fp", [TypedAttrInterface]> { // ConstArrayAttr //===----------------------------------------------------------------------===// -def ConstArrayAttr : CIR_Attr<"ConstArray", "const_array", [TypedAttrInterface]> { +def CIR_ConstArrayAttr : CIR_Attr<"ConstArray", "const_array", [ + TypedAttrInterface +]> { let summary = "A constant array from ArrayAttr or StringRefAttr"; let description = [{ An CIR array attribute is an array of literals of the specified attr types. @@ -310,8 +312,9 @@ def ConstArrayAttr : CIR_Attr<"ConstArray", "const_array", [TypedAttrInterface]> // ConstVectorAttr //===----------------------------------------------------------------------===// -def ConstVectorAttr : CIR_Attr<"ConstVector", "const_vector", - [TypedAttrInterface]> { +def CIR_ConstVectorAttr : CIR_Attr<"ConstVector", "const_vector", [ + TypedAttrInterface +]> { let summary = "A constant vector from ArrayAttr"; let description = [{ A CIR vector attribute is an array of literals of the specified attribute @@ -342,7 +345,7 @@ def ConstVectorAttr : CIR_Attr<"ConstVector", "const_vector", // ConstPtrAttr //===----------------------------------------------------------------------===// -def ConstPtrAttr : CIR_Attr<"ConstPtr", "ptr", [TypedAttrInterface]> { +def CIR_ConstPtrAttr : CIR_Attr<"ConstPtr", "ptr", [TypedAttrInterface]> { let summary = "Holds a constant pointer value"; let parameters = (ins AttributeSelfTypeParameter<"", "::cir::PointerType">:$type, @@ -371,8 +374,9 @@ def ConstPtrAttr : CIR_Attr<"ConstPtr", "ptr", [TypedAttrInterface]> { // ConstComplexAttr //===----------------------------------------------------------------------===// -def ConstComplexAttr : CIR_Attr<"ConstComplex", "const_complex", - [TypedAttrInterface]> { +def CIR_ConstComplexAttr : CIR_Attr<"ConstComplex", "const_complex", [ + TypedAttrInterface +]> { let summary = "An attribute that contains a constant complex value"; let description = [{ The `#cir.const_complex` attribute contains a constant value of complex @@ -454,7 +458,7 @@ def CIR_VisibilityAttr : CIR_EnumAttr<CIR_VisibilityKind, "visibility"> { // BitfieldInfoAttr //===----------------------------------------------------------------------===// -def BitfieldInfoAttr : CIR_Attr<"BitfieldInfo", "bitfield_info"> { +def CIR_BitfieldInfoAttr : CIR_Attr<"BitfieldInfo", "bitfield_info"> { let summary = "Represents info for a bit-field member"; let description = [{ Holds the following information about bitfields: name, storage type, size @@ -512,5 +516,4 @@ def BitfieldInfoAttr : CIR_Attr<"BitfieldInfo", "bitfield_info"> { ]; } - #endif // CLANG_CIR_DIALECT_IR_CIRATTRS_TD diff --git a/clang/include/clang/CIR/Dialect/IR/CIROps.td b/clang/include/clang/CIR/Dialect/IR/CIROps.td index e2ddbd1..5ef5b60 100644 --- a/clang/include/clang/CIR/Dialect/IR/CIROps.td +++ b/clang/include/clang/CIR/Dialect/IR/CIROps.td @@ -607,8 +607,8 @@ def CIR_ConditionOp : CIR_Op<"condition", [ //===----------------------------------------------------------------------===// defvar CIR_YieldableScopes = [ - "ArrayCtor", "CaseOp", "DoWhileOp", "ForOp", "IfOp", "ScopeOp", "SwitchOp", - "TernaryOp", "WhileOp" + "ArrayCtor", "ArrayDtor", "CaseOp", "DoWhileOp", "ForOp", "IfOp", "ScopeOp", + "SwitchOp", "TernaryOp", "WhileOp" ]; def CIR_YieldOp : CIR_Op<"yield", [ @@ -1747,7 +1747,7 @@ def CIR_SetBitfieldOp : CIR_Op<"set_bitfield"> { let arguments = (ins Arg<CIR_PointerType, "the address to store the value", [MemWrite]>:$addr, CIR_AnyType:$src, - BitfieldInfoAttr:$bitfield_info, + CIR_BitfieldInfoAttr:$bitfield_info, DefaultValuedOptionalAttr<I64Attr, "0">:$alignment, UnitAttr:$is_volatile ); @@ -1834,7 +1834,7 @@ def CIR_GetBitfieldOp : CIR_Op<"get_bitfield"> { let arguments = (ins Arg<CIR_PointerType, "the address to load from", [MemRead]>:$addr, - BitfieldInfoAttr:$bitfield_info, + CIR_BitfieldInfoAttr:$bitfield_info, DefaultValuedOptionalAttr<I64Attr, "0">:$alignment, UnitAttr:$is_volatile ); @@ -1946,6 +1946,10 @@ def CIR_FuncOp : CIR_Op<"func", [ The function linkage information is specified by `linkage`, as defined by `GlobalLinkageKind` attribute. + The `no_proto` keyword is used to identify functions that were declared + without a prototype and, consequently, may contain calls with invalid + arguments and undefined behavior. + Example: ```mlir @@ -1964,6 +1968,7 @@ def CIR_FuncOp : CIR_Op<"func", [ let arguments = (ins SymbolNameAttr:$sym_name, CIR_VisibilityAttr:$global_visibility, TypeAttrOf<CIR_FuncType>:$function_type, + UnitAttr:$no_proto, UnitAttr:$dso_local, DefaultValuedAttr<CIR_GlobalLinkageKind, "cir::GlobalLinkageKind::ExternalLinkage">:$linkage, @@ -2005,13 +2010,6 @@ def CIR_FuncOp : CIR_Op<"func", [ return getFunctionType().getReturnTypes(); } - // TODO(cir): this should be an operand attribute, but for now we just hard- - // wire this as a function. Will later add a $no_proto argument to this op. - bool getNoProto() { - assert(!cir::MissingFeatures::opFuncNoProto()); - return false; - } - //===------------------------------------------------------------------===// // SymbolOpInterface Methods //===------------------------------------------------------------------===// @@ -2229,7 +2227,7 @@ def CIR_TrapOp : CIR_Op<"trap", [Terminator]> { } //===----------------------------------------------------------------------===// -// ArrayCtor +// ArrayCtor & ArrayDtor //===----------------------------------------------------------------------===// class CIR_ArrayInitDestroy<string mnemonic> : CIR_Op<mnemonic> { @@ -2260,7 +2258,9 @@ def CIR_ArrayCtor : CIR_ArrayInitDestroy<"array.ctor"> { let description = [{ Initialize each array element using the same C++ constructor. This operation has one region, with one single block. The block has an - incoming argument for the current array index to initialize. + incoming argument for the current array element to initialize. + + Example: ```mlir cir.array.ctor(%0 : !cir.ptr<!cir.array<!rec_S x 42>>) { @@ -2272,6 +2272,25 @@ def CIR_ArrayCtor : CIR_ArrayInitDestroy<"array.ctor"> { }]; } +def CIR_ArrayDtor : CIR_ArrayInitDestroy<"array.dtor"> { + let summary = "Destroy array elements with C++ dtors"; + let description = [{ + Destroy each array element using the same C++ destructor. This + operation has one region, with one single block. The block has an + incoming argument for the current array element to destruct. + + Example: + + ```mlir + cir.array.dtor(%0 : !cir.ptr<!cir.array<!rec_S x 42>>) { + ^bb0(%arg0: !cir.ptr<!rec_S>): + cir.call @some_dtor(%arg0) : (!cir.ptr<!rec_S>) -> () + cir.yield + } + ``` + }]; +} + //===----------------------------------------------------------------------===// // VecCreate //===----------------------------------------------------------------------===// @@ -2804,6 +2823,53 @@ def CIR_ComplexSubOp : CIR_Op<"complex.sub", [ }]; } +//===----------------------------------------------------------------------===// +// ComplexMulOp +//===----------------------------------------------------------------------===// + +def CIR_ComplexRangeKind : CIR_I32EnumAttr< + "ComplexRangeKind", "complex multiplication and division implementation", [ + I32EnumAttrCase<"Full", 0, "full">, + I32EnumAttrCase<"Improved", 1, "improved">, + I32EnumAttrCase<"Promoted", 2, "promoted">, + I32EnumAttrCase<"Basic", 3, "basic">, +]>; + +def CIR_ComplexMulOp : CIR_Op<"complex.mul", [ + Pure, SameOperandsAndResultType +]> { + let summary = "Complex multiplication"; + let description = [{ + The `cir.complex.mul` operation takes two complex numbers and returns + their product. + + Range is used to select the implementation used when the operation + is lowered to the LLVM dialect. For multiplication, 'improved', + 'promoted', and 'basic' are all handled equivalently, producing the + algebraic formula with no special handling for NaN value. If 'full' is + used, a runtime-library function is called if one of the intermediate + calculations produced a NaN value. + + Example: + + ```mlir + %2 = cir.complex.mul %0, %1 range(basic) : !cir.complex<!cir.float> + %2 = cir.complex.mul %0, %1 range(full) : !cir.complex<!cir.float> + ``` + }]; + + let arguments = (ins + CIR_ComplexType:$lhs, + CIR_ComplexType:$rhs, + CIR_ComplexRangeKind:$range + ); + + let results = (outs CIR_ComplexType:$result); + + let assemblyFormat = [{ + $lhs `,` $rhs `range` `(` $range `)` `:` qualified(type($result)) attr-dict + }]; +} //===----------------------------------------------------------------------===// // Bit Manipulation Operations @@ -2911,6 +2977,28 @@ def CIR_BitCtzOp : CIR_BitZeroCountOpBase<"ctz", }]; } +def CIR_BitFfsOp : CIR_BitOpBase<"ffs", CIR_SIntOfWidths<[32, 64]>> { + let summary = "Get the position of the least significant 1-bit in input"; + let description = [{ + Compute the 1-based position of the least significant 1-bit of the input. + + The input integer must be a signed integer. The `cir.ffs` operation returns + one plus the index of the least significant 1-bit of the input signed + integer. If the input integer is 0, `cir.ffs` yields 0. + + Example: + + ```mlir + !s32i = !cir.int<s, 32> + + // %0 = 0x0010_1000 + %0 = cir.const #cir.int<40> : !s32i + // #1 will be 4 since the 4th least significant bit is 1. + %1 = cir.ffs %0 : !s32i + ``` + }]; +} + def CIR_BitParityOp : CIR_BitOpBase<"parity", CIR_UIntOfWidths<[32, 64]>> { let summary = "Get the parity of input"; let description = [{ diff --git a/clang/include/clang/CIR/MissingFeatures.h b/clang/include/clang/CIR/MissingFeatures.h index e1a5c3d..adc7b5f 100644 --- a/clang/include/clang/CIR/MissingFeatures.h +++ b/clang/include/clang/CIR/MissingFeatures.h @@ -73,14 +73,16 @@ struct MissingFeatures { // FuncOp handling static bool opFuncOpenCLKernelMetadata() { return false; } static bool opFuncAstDeclAttr() { return false; } + static bool opFuncAttributesForDefinition() { return false; } static bool opFuncCallingConv() { return false; } - static bool opFuncExtraAttrs() { return false; } - static bool opFuncNoProto() { return false; } static bool opFuncCPUAndFeaturesAttributes() { return false; } - static bool opFuncSection() { return false; } - static bool opFuncMultipleReturnVals() { return false; } - static bool opFuncAttributesForDefinition() { return false; } + static bool opFuncExceptions() { return false; } + static bool opFuncExtraAttrs() { return false; } static bool opFuncMaybeHandleStaticInExternC() { return false; } + static bool opFuncMultipleReturnVals() { return false; } + static bool opFuncOperandBundles() { return false; } + static bool opFuncParameterAttributes() { return false; } + static bool opFuncSection() { return false; } static bool setLLVMFunctionFEnvAttributes() { return false; } static bool setFunctionAttributes() { return false; } @@ -96,7 +98,6 @@ struct MissingFeatures { static bool opCallReturn() { return false; } static bool opCallArgEvaluationOrder() { return false; } static bool opCallCallConv() { return false; } - static bool opCallNoPrototypeFunc() { return false; } static bool opCallMustTail() { return false; } static bool opCallVirtual() { return false; } static bool opCallInAlloca() { return false; } @@ -109,6 +110,7 @@ struct MissingFeatures { static bool opCallCIRGenFuncInfoExtParamInfo() { return false; } static bool opCallLandingPad() { return false; } static bool opCallContinueBlock() { return false; } + static bool opCallChain() { return false; } // CXXNewExpr static bool exprNewNullCheck() { return false; } @@ -215,6 +217,7 @@ struct MissingFeatures { static bool intrinsics() { return false; } static bool isMemcpyEquivalentSpecialMember() { return false; } static bool isTrivialCtorOrDtor() { return false; } + static bool lambdaCaptures() { return false; } static bool lambdaFieldToName() { return false; } static bool loopInfoStack() { return false; } static bool lowerAggregateLoadStore() { return false; } @@ -224,7 +227,6 @@ struct MissingFeatures { static bool moduleNameHash() { return false; } static bool msabi() { return false; } static bool needsGlobalCtorDtor() { return false; } - static bool nonFineGrainedBitfields() { return false; } static bool objCBlocks() { return false; } static bool objCGC() { return false; } static bool objCLifetime() { return false; } diff --git a/clang/include/clang/Driver/CudaInstallationDetector.h b/clang/include/clang/Driver/CudaInstallationDetector.h new file mode 100644 index 0000000..0fecbfe --- /dev/null +++ b/clang/include/clang/Driver/CudaInstallationDetector.h @@ -0,0 +1,76 @@ +//===-- CudaInstallationDetector.h - Cuda Instalation Detector --*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_DRIVER_CUDAINSTALLATIONDETECTOR_H +#define LLVM_CLANG_DRIVER_CUDAINSTALLATIONDETECTOR_H + +#include "clang/Basic/Cuda.h" +#include "clang/Driver/Driver.h" +#include <bitset> + +namespace clang { +namespace driver { + +/// A class to find a viable CUDA installation +class CudaInstallationDetector { +private: + const Driver &D; + bool IsValid = false; + CudaVersion Version = CudaVersion::UNKNOWN; + std::string InstallPath; + std::string BinPath; + std::string LibDevicePath; + std::string IncludePath; + llvm::StringMap<std::string> LibDeviceMap; + + // CUDA architectures for which we have raised an error in + // CheckCudaVersionSupportsArch. + mutable std::bitset<(int)OffloadArch::LAST> ArchsWithBadVersion; + +public: + CudaInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, + const llvm::opt::ArgList &Args); + + void AddCudaIncludeArgs(const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args) const; + + /// Emit an error if Version does not support the given Arch. + /// + /// If either Version or Arch is unknown, does not emit an error. Emits at + /// most one error per Arch. + void CheckCudaVersionSupportsArch(OffloadArch Arch) const; + + /// Check whether we detected a valid Cuda install. + bool isValid() const { return IsValid; } + /// Print information about the detected CUDA installation. + void print(raw_ostream &OS) const; + + /// Get the detected Cuda install's version. + CudaVersion version() const { + return Version == CudaVersion::NEW ? CudaVersion::PARTIALLY_SUPPORTED + : Version; + } + /// Get the detected Cuda installation path. + StringRef getInstallPath() const { return InstallPath; } + /// Get the detected path to Cuda's bin directory. + StringRef getBinPath() const { return BinPath; } + /// Get the detected Cuda Include path. + StringRef getIncludePath() const { return IncludePath; } + /// Get the detected Cuda device library path. + StringRef getLibDevicePath() const { return LibDevicePath; } + /// Get libdevice file for given architecture + std::string getLibDeviceFile(StringRef Gpu) const { + return LibDeviceMap.lookup(Gpu); + } + void WarnIfUnsupportedVersion() const; +}; + +} // namespace driver +} // namespace clang + +#endif // LLVM_CLANG_DRIVER_CUDAINSTALLATIONDETECTOR_H diff --git a/clang/include/clang/Driver/LazyDetector.h b/clang/include/clang/Driver/LazyDetector.h new file mode 100644 index 0000000..5bd1d01 --- /dev/null +++ b/clang/include/clang/Driver/LazyDetector.h @@ -0,0 +1,45 @@ +//===--- LazyDetector.h - Lazy ToolChain Detection --------------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_DRIVER_LAZYDETECTOR_H +#define LLVM_CLANG_DRIVER_LAZYDETECTOR_H + +#include "clang/Driver/Tool.h" +#include "clang/Driver/ToolChain.h" +#include <optional> + +namespace clang { + +/// Simple wrapper for toolchain detector with costly initialization. This +/// delays the creation of the actual detector until its first usage. + +template <class T> class LazyDetector { + const driver::Driver &D; + llvm::Triple Triple; + const llvm::opt::ArgList &Args; + + std::optional<T> Detector; + +public: + LazyDetector(const driver::Driver &D, const llvm::Triple &Triple, + const llvm::opt::ArgList &Args) + : D(D), Triple(Triple), Args(Args) {} + T *operator->() { + if (!Detector) + Detector.emplace(D, Triple, Args); + return &*Detector; + } + const T *operator->() const { + return const_cast<T const *>( + const_cast<LazyDetector &>(*this).operator->()); + } +}; + +} // end namespace clang + +#endif // LLVM_CLANG_DRIVER_LAZYDETECTOR_H diff --git a/clang/include/clang/Driver/Options.td b/clang/include/clang/Driver/Options.td index eb53821..3c04aeb 100644 --- a/clang/include/clang/Driver/Options.td +++ b/clang/include/clang/Driver/Options.td @@ -5312,6 +5312,8 @@ def mextended_const : Flag<["-"], "mextended-const">, Group<m_wasm_Features_Grou def mno_extended_const : Flag<["-"], "mno-extended-const">, Group<m_wasm_Features_Group>; def mfp16 : Flag<["-"], "mfp16">, Group<m_wasm_Features_Group>; def mno_fp16 : Flag<["-"], "mno-fp16">, Group<m_wasm_Features_Group>; +def mgc : Flag<["-"], "mgc">, Group<m_wasm_Features_Group>; +def mno_gc : Flag<["-"], "mno-gc">, Group<m_wasm_Features_Group>; def mmultimemory : Flag<["-"], "mmultimemory">, Group<m_wasm_Features_Group>; def mno_multimemory : Flag<["-"], "mno-multimemory">, Group<m_wasm_Features_Group>; def mmultivalue : Flag<["-"], "mmultivalue">, Group<m_wasm_Features_Group>; diff --git a/clang/include/clang/Driver/RocmInstallationDetector.h b/clang/include/clang/Driver/RocmInstallationDetector.h new file mode 100644 index 0000000..3a325a6 --- /dev/null +++ b/clang/include/clang/Driver/RocmInstallationDetector.h @@ -0,0 +1,303 @@ +//===-- RocmInstallationDetector.h - ROCm Instalation Detector --*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_DRIVER_ROCMINSTALLATIONDETECTOR_H +#define LLVM_CLANG_DRIVER_ROCMINSTALLATIONDETECTOR_H + +#include "clang/Driver/Driver.h" + +namespace clang { +namespace driver { + +/// ABI version of device library. +struct DeviceLibABIVersion { + unsigned ABIVersion = 0; + DeviceLibABIVersion(unsigned V) : ABIVersion(V) {} + static DeviceLibABIVersion fromCodeObjectVersion(unsigned CodeObjectVersion) { + if (CodeObjectVersion < 4) + CodeObjectVersion = 4; + return DeviceLibABIVersion(CodeObjectVersion * 100); + } + /// Whether ABI version bc file is requested. + /// ABIVersion is code object version multiplied by 100. Code object v4 + /// and below works with ROCm 5.0 and below which does not have + /// abi_version_*.bc. Code object v5 requires abi_version_500.bc. + bool requiresLibrary() { return ABIVersion >= 500; } + std::string toString() { return Twine(getAsCodeObjectVersion()).str(); } + + unsigned getAsCodeObjectVersion() const { + assert(ABIVersion % 100 == 0 && "Not supported"); + return ABIVersion / 100; + } +}; + +/// A class to find a viable ROCM installation +/// TODO: Generalize to handle libclc. +class RocmInstallationDetector { +private: + struct ConditionalLibrary { + SmallString<0> On; + SmallString<0> Off; + + bool isValid() const { return !On.empty() && !Off.empty(); } + + StringRef get(bool Enabled) const { + assert(isValid()); + return Enabled ? On : Off; + } + }; + + // Installation path candidate. + struct Candidate { + llvm::SmallString<0> Path; + bool StrictChecking; + // Release string for ROCm packages built with SPACK if not empty. The + // installation directories of ROCm packages built with SPACK follow the + // convention <package_name>-<rocm_release_string>-<hash>. + std::string SPACKReleaseStr; + + bool isSPACK() const { return !SPACKReleaseStr.empty(); } + Candidate(std::string Path, bool StrictChecking = false, + StringRef SPACKReleaseStr = {}) + : Path(Path), StrictChecking(StrictChecking), + SPACKReleaseStr(SPACKReleaseStr.str()) {} + }; + + struct CommonBitcodeLibsPreferences { + CommonBitcodeLibsPreferences(const Driver &D, + const llvm::opt::ArgList &DriverArgs, + StringRef GPUArch, + const Action::OffloadKind DeviceOffloadingKind, + const bool NeedsASanRT); + + DeviceLibABIVersion ABIVer; + bool IsOpenMP; + bool Wave64; + bool DAZ; + bool FiniteOnly; + bool UnsafeMathOpt; + bool FastRelaxedMath; + bool CorrectSqrt; + bool GPUSan; + }; + + const Driver &D; + bool HasHIPRuntime = false; + bool HasDeviceLibrary = false; + bool HasHIPStdParLibrary = false; + bool HasRocThrustLibrary = false; + bool HasRocPrimLibrary = false; + + // Default version if not detected or specified. + const unsigned DefaultVersionMajor = 3; + const unsigned DefaultVersionMinor = 5; + const char *DefaultVersionPatch = "0"; + + // The version string in Major.Minor.Patch format. + std::string DetectedVersion; + // Version containing major and minor. + llvm::VersionTuple VersionMajorMinor; + // Version containing patch. + std::string VersionPatch; + + // ROCm path specified by --rocm-path. + StringRef RocmPathArg; + // ROCm device library paths specified by --rocm-device-lib-path. + std::vector<std::string> RocmDeviceLibPathArg; + // HIP runtime path specified by --hip-path. + StringRef HIPPathArg; + // HIP Standard Parallel Algorithm acceleration library specified by + // --hipstdpar-path + StringRef HIPStdParPathArg; + // rocThrust algorithm library specified by --hipstdpar-thrust-path + StringRef HIPRocThrustPathArg; + // rocPrim algorithm library specified by --hipstdpar-prim-path + StringRef HIPRocPrimPathArg; + // HIP version specified by --hip-version. + StringRef HIPVersionArg; + // Wheter -nogpulib is specified. + bool NoBuiltinLibs = false; + + // Paths + SmallString<0> InstallPath; + SmallString<0> BinPath; + SmallString<0> LibPath; + SmallString<0> LibDevicePath; + SmallString<0> IncludePath; + SmallString<0> SharePath; + llvm::StringMap<std::string> LibDeviceMap; + + // Libraries that are always linked. + SmallString<0> OCML; + SmallString<0> OCKL; + + // Libraries that are always linked depending on the language + SmallString<0> OpenCL; + + // Asan runtime library + SmallString<0> AsanRTL; + + // Libraries swapped based on compile flags. + ConditionalLibrary WavefrontSize64; + ConditionalLibrary FiniteOnly; + ConditionalLibrary UnsafeMath; + ConditionalLibrary DenormalsAreZero; + ConditionalLibrary CorrectlyRoundedSqrt; + + // Maps ABI version to library path. The version number is in the format of + // three digits as used in the ABI version library name. + std::map<unsigned, std::string> ABIVersionMap; + + // Cache ROCm installation search paths. + SmallVector<Candidate, 4> ROCmSearchDirs; + bool PrintROCmSearchDirs; + bool Verbose; + + bool allGenericLibsValid() const { + return !OCML.empty() && !OCKL.empty() && !OpenCL.empty() && + WavefrontSize64.isValid() && FiniteOnly.isValid() && + UnsafeMath.isValid() && DenormalsAreZero.isValid() && + CorrectlyRoundedSqrt.isValid(); + } + + void scanLibDevicePath(llvm::StringRef Path); + bool parseHIPVersionFile(llvm::StringRef V); + const SmallVectorImpl<Candidate> &getInstallationPathCandidates(); + + /// Find the path to a SPACK package under the ROCm candidate installation + /// directory if the candidate is a SPACK ROCm candidate. \returns empty + /// string if the candidate is not SPACK ROCm candidate or the requested + /// package is not found. + llvm::SmallString<0> findSPACKPackage(const Candidate &Cand, + StringRef PackageName); + +public: + RocmInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, + const llvm::opt::ArgList &Args, + bool DetectHIPRuntime = true, + bool DetectDeviceLib = false); + + /// Get file paths of default bitcode libraries common to AMDGPU based + /// toolchains. + llvm::SmallVector<ToolChain::BitCodeLibraryInfo, 12> + getCommonBitcodeLibs(const llvm::opt::ArgList &DriverArgs, + StringRef LibDeviceFile, StringRef GPUArch, + const Action::OffloadKind DeviceOffloadingKind, + const bool NeedsASanRT) const; + /// Check file paths of default bitcode libraries common to AMDGPU based + /// toolchains. \returns false if there are invalid or missing files. + bool checkCommonBitcodeLibs(StringRef GPUArch, StringRef LibDeviceFile, + DeviceLibABIVersion ABIVer) const; + + /// Check whether we detected a valid HIP runtime. + bool hasHIPRuntime() const { return HasHIPRuntime; } + + /// Check whether we detected a valid ROCm device library. + bool hasDeviceLibrary() const { return HasDeviceLibrary; } + + /// Check whether we detected a valid HIP STDPAR Acceleration library. + bool hasHIPStdParLibrary() const { return HasHIPStdParLibrary; } + + /// Print information about the detected ROCm installation. + void print(raw_ostream &OS) const; + + /// Get the detected Rocm install's version. + // RocmVersion version() const { return Version; } + + /// Get the detected Rocm installation path. + StringRef getInstallPath() const { return InstallPath; } + + /// Get the detected path to Rocm's bin directory. + // StringRef getBinPath() const { return BinPath; } + + /// Get the detected Rocm Include path. + StringRef getIncludePath() const { return IncludePath; } + + /// Get the detected Rocm library path. + StringRef getLibPath() const { return LibPath; } + + /// Get the detected Rocm device library path. + StringRef getLibDevicePath() const { return LibDevicePath; } + + StringRef getOCMLPath() const { + assert(!OCML.empty()); + return OCML; + } + + StringRef getOCKLPath() const { + assert(!OCKL.empty()); + return OCKL; + } + + StringRef getOpenCLPath() const { + assert(!OpenCL.empty()); + return OpenCL; + } + + /// Returns empty string of Asan runtime library is not available. + StringRef getAsanRTLPath() const { return AsanRTL; } + + StringRef getWavefrontSize64Path(bool Enabled) const { + return WavefrontSize64.get(Enabled); + } + + StringRef getFiniteOnlyPath(bool Enabled) const { + return FiniteOnly.get(Enabled); + } + + StringRef getUnsafeMathPath(bool Enabled) const { + return UnsafeMath.get(Enabled); + } + + StringRef getDenormalsAreZeroPath(bool Enabled) const { + return DenormalsAreZero.get(Enabled); + } + + StringRef getCorrectlyRoundedSqrtPath(bool Enabled) const { + return CorrectlyRoundedSqrt.get(Enabled); + } + + StringRef getABIVersionPath(DeviceLibABIVersion ABIVer) const { + auto Loc = ABIVersionMap.find(ABIVer.ABIVersion); + if (Loc == ABIVersionMap.end()) + return StringRef(); + return Loc->second; + } + + /// Get libdevice file for given architecture + StringRef getLibDeviceFile(StringRef Gpu) const { + auto Loc = LibDeviceMap.find(Gpu); + if (Loc == LibDeviceMap.end()) + return ""; + return Loc->second; + } + + void AddHIPIncludeArgs(const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args) const; + + void detectDeviceLibrary(); + void detectHIPRuntime(); + + /// Get the values for --rocm-device-lib-path arguments + ArrayRef<std::string> getRocmDeviceLibPathArg() const { + return RocmDeviceLibPathArg; + } + + /// Get the value for --rocm-path argument + StringRef getRocmPathArg() const { return RocmPathArg; } + + /// Get the value for --hip-version argument + StringRef getHIPVersionArg() const { return HIPVersionArg; } + + StringRef getHIPVersion() const { return DetectedVersion; } +}; + +} // namespace driver +} // namespace clang + +#endif // LLVM_CLANG_DRIVER_ROCMINSTALLATIONDETECTOR_H diff --git a/clang/include/clang/Driver/SyclInstallationDetector.h b/clang/include/clang/Driver/SyclInstallationDetector.h new file mode 100644 index 0000000..6925ec2 --- /dev/null +++ b/clang/include/clang/Driver/SyclInstallationDetector.h @@ -0,0 +1,29 @@ +//===-- SyclInstallationDetector.h - SYCL Instalation Detector --*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef LLVM_CLANG_DRIVER_SYCLINSTALLATIONDETECTOR_H +#define LLVM_CLANG_DRIVER_SYCLINSTALLATIONDETECTOR_H + +#include "clang/Driver/Driver.h" + +namespace clang { +namespace driver { + +class SYCLInstallationDetector { +public: + SYCLInstallationDetector(const Driver &D, const llvm::Triple &HostTriple, + const llvm::opt::ArgList &Args); + + void addSYCLIncludeArgs(const llvm::opt::ArgList &DriverArgs, + llvm::opt::ArgStringList &CC1Args) const; +}; + +} // namespace driver +} // namespace clang + +#endif // LLVM_CLANG_DRIVER_SYCLINSTALLATIONDETECTOR_H diff --git a/clang/include/clang/Frontend/ASTUnit.h b/clang/include/clang/Frontend/ASTUnit.h index 1286fe4..7dd9aef 100644 --- a/clang/include/clang/Frontend/ASTUnit.h +++ b/clang/include/clang/Frontend/ASTUnit.h @@ -445,6 +445,9 @@ public: const DiagnosticsEngine &getDiagnostics() const { return *Diagnostics; } DiagnosticsEngine &getDiagnostics() { return *Diagnostics; } + llvm::IntrusiveRefCntPtr<DiagnosticsEngine> getDiagnosticsPtr() { + return Diagnostics; + } const SourceManager &getSourceManager() const { return *SourceMgr; } SourceManager &getSourceManager() { return *SourceMgr; } @@ -918,8 +921,9 @@ public: bool IncludeCodePatterns, bool IncludeBriefComments, CodeCompleteConsumer &Consumer, std::shared_ptr<PCHContainerOperations> PCHContainerOps, - DiagnosticsEngine &Diag, LangOptions &LangOpts, - SourceManager &SourceMgr, FileManager &FileMgr, + llvm::IntrusiveRefCntPtr<DiagnosticsEngine> Diag, + LangOptions &LangOpts, SourceManager &SourceMgr, + FileManager &FileMgr, SmallVectorImpl<StoredDiagnostic> &StoredDiagnostics, SmallVectorImpl<const llvm::MemoryBuffer *> &OwnedBuffers, std::unique_ptr<SyntaxOnlyAction> Act); diff --git a/clang/include/clang/Frontend/CompilerInstance.h b/clang/include/clang/Frontend/CompilerInstance.h index 2408367..a24decd 100644 --- a/clang/include/clang/Frontend/CompilerInstance.h +++ b/clang/include/clang/Frontend/CompilerInstance.h @@ -361,7 +361,7 @@ public: } /// setDiagnostics - Replace the current diagnostics engine. - void setDiagnostics(DiagnosticsEngine *Value); + void setDiagnostics(llvm::IntrusiveRefCntPtr<DiagnosticsEngine> Value); DiagnosticConsumer &getDiagnosticClient() const { assert(Diagnostics && Diagnostics->getClient() && @@ -420,6 +420,8 @@ public: /// @{ llvm::vfs::FileSystem &getVirtualFileSystem() const; + llvm::IntrusiveRefCntPtr<llvm::vfs::FileSystem> + getVirtualFileSystemPtr() const; /// @} /// @name File Manager diff --git a/clang/include/clang/Frontend/FrontendAction.h b/clang/include/clang/Frontend/FrontendAction.h index 718684a..08c5fbc 100644 --- a/clang/include/clang/Frontend/FrontendAction.h +++ b/clang/include/clang/Frontend/FrontendAction.h @@ -84,6 +84,8 @@ protected: /// \return True on success; on failure ExecutionAction() and /// EndSourceFileAction() will not be called. virtual bool BeginSourceFileAction(CompilerInstance &CI) { + if (CurrentInput.isPreprocessed()) + CI.getPreprocessor().SetMacroExpansionOnlyInDirectives(); return true; } @@ -98,7 +100,11 @@ protected: /// /// This is guaranteed to only be called following a successful call to /// BeginSourceFileAction (and BeginSourceFile). - virtual void EndSourceFileAction() {} + virtual void EndSourceFileAction() { + if (CurrentInput.isPreprocessed()) + // Reset the preprocessor macro expansion to the default. + getCompilerInstance().getPreprocessor().SetEnableMacroExpansion(); + } /// Callback at the end of processing a single input, to determine /// if the output files should be erased or not. diff --git a/clang/include/clang/Frontend/PrecompiledPreamble.h b/clang/include/clang/Frontend/PrecompiledPreamble.h index 624df00..565395b 100644 --- a/clang/include/clang/Frontend/PrecompiledPreamble.h +++ b/clang/include/clang/Frontend/PrecompiledPreamble.h @@ -84,7 +84,7 @@ public: static llvm::ErrorOr<PrecompiledPreamble> Build(const CompilerInvocation &Invocation, const llvm::MemoryBuffer *MainFileBuffer, PreambleBounds Bounds, - DiagnosticsEngine &Diagnostics, + IntrusiveRefCntPtr<DiagnosticsEngine> Diagnostics, IntrusiveRefCntPtr<llvm::vfs::FileSystem> VFS, std::shared_ptr<PCHContainerOperations> PCHContainerOps, bool StoreInMemory, StringRef StoragePath, diff --git a/clang/include/clang/Lex/Preprocessor.h b/clang/include/clang/Lex/Preprocessor.h index 4d82e20..71b0f8e 100644 --- a/clang/include/clang/Lex/Preprocessor.h +++ b/clang/include/clang/Lex/Preprocessor.h @@ -1847,6 +1847,10 @@ public: MacroExpansionInDirectivesOverride = true; } + void SetEnableMacroExpansion() { + DisableMacroExpansion = MacroExpansionInDirectivesOverride = false; + } + /// Peeks ahead N tokens and returns that token without consuming any /// tokens. /// diff --git a/clang/include/clang/Sema/SemaARM.h b/clang/include/clang/Sema/SemaARM.h index e77d65f..104992e 100644 --- a/clang/include/clang/Sema/SemaARM.h +++ b/clang/include/clang/Sema/SemaARM.h @@ -44,8 +44,8 @@ public: bool CheckImmediateArg(CallExpr *TheCall, unsigned CheckTy, unsigned ArgIdx, unsigned EltBitWidth, unsigned VecBitWidth); - bool CheckARMBuiltinExclusiveCall(unsigned BuiltinID, CallExpr *TheCall, - unsigned MaxWidth); + bool CheckARMBuiltinExclusiveCall(const TargetInfo &TI, unsigned BuiltinID, + CallExpr *TheCall); bool CheckNeonBuiltinFunctionCall(const TargetInfo &TI, unsigned BuiltinID, CallExpr *TheCall); bool PerformNeonImmChecks( diff --git a/clang/include/clang/StaticAnalyzer/Checkers/Checkers.td b/clang/include/clang/StaticAnalyzer/Checkers/Checkers.td index 6225977..36196cd 100644 --- a/clang/include/clang/StaticAnalyzer/Checkers/Checkers.td +++ b/clang/include/clang/StaticAnalyzer/Checkers/Checkers.td @@ -233,15 +233,11 @@ def UndefResultChecker : Checker<"UndefinedBinaryOperatorResult">, HelpText<"Check for undefined results of binary operators">, Documentation<HasDocumentation>; -def StackAddrEscapeBase : Checker<"StackAddrEscapeBase">, - HelpText<"Generate information about stack address escapes.">, - Documentation<NotDocumented>, - Hidden; - -def StackAddrEscapeChecker : Checker<"StackAddressEscape">, - HelpText<"Check that addresses to stack memory do not escape the function">, - Dependencies<[StackAddrEscapeBase]>, - Documentation<HasDocumentation>; +def StackAddrEscapeChecker + : Checker<"StackAddressEscape">, + HelpText< + "Check that addresses to stack memory do not escape the function">, + Documentation<HasDocumentation>; def DynamicTypePropagation : Checker<"DynamicTypePropagation">, HelpText<"Generate dynamic type information">, @@ -295,10 +291,11 @@ def DynamicTypeChecker Dependencies<[DynamicTypePropagation]>, Documentation<HasDocumentation>; -def StackAddrAsyncEscapeChecker : Checker<"StackAddressAsyncEscape">, - HelpText<"Check that addresses to stack memory do not escape the function">, - Dependencies<[StackAddrEscapeBase]>, - Documentation<HasDocumentation>; +def StackAddrAsyncEscapeChecker + : Checker<"StackAddressAsyncEscape">, + HelpText< + "Check that addresses to stack memory do not escape the function">, + Documentation<HasDocumentation>; def PthreadLockBase : Checker<"PthreadLockBase">, HelpText<"Helper registering multiple checks.">, @@ -1047,11 +1044,6 @@ def RetainCountBase : Checker<"RetainCountBase">, let ParentPackage = OSX in { -def NSOrCFErrorDerefChecker : Checker<"NSOrCFErrorDerefChecker">, - HelpText<"Implementation checker for NSErrorChecker and CFErrorChecker">, - Documentation<NotDocumented>, - Hidden; - def NumberObjectConversionChecker : Checker<"NumberObjectConversion">, HelpText<"Check for erroneous conversions of objects representing numbers " "into numbers">, @@ -1149,9 +1141,8 @@ def ObjCSuperCallChecker : Checker<"MissingSuperCall">, Documentation<NotDocumented>; def NSErrorChecker : Checker<"NSError">, - HelpText<"Check usage of NSError** parameters">, - Dependencies<[NSOrCFErrorDerefChecker]>, - Documentation<HasDocumentation>; + HelpText<"Check usage of NSError** parameters">, + Documentation<HasDocumentation>; def RetainCountChecker : Checker<"RetainCount">, HelpText<"Check for leaks and improper reference count management">, @@ -1253,9 +1244,8 @@ def CFRetainReleaseChecker : Checker<"CFRetainRelease">, Documentation<HasDocumentation>; def CFErrorChecker : Checker<"CFError">, - HelpText<"Check usage of CFErrorRef* parameters">, - Dependencies<[NSOrCFErrorDerefChecker]>, - Documentation<HasDocumentation>; + HelpText<"Check usage of CFErrorRef* parameters">, + Documentation<HasDocumentation>; } // end "osx.coreFoundation" |