llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang-modules @llvm/pr-subscribers-clang @llvm/pr-subscribers-debuginfo Author: Shilei Tian (shiltian) <details> <summary>Changes</summary> - **[Clang][AMDGPU] Add a new builtin type for buffer rsrc** - **[Clang][AMDGPU] Add a builtin for `llvm.amdgcn.make.buffer.rsrc` intrinsic** --- Patch is 28.13 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/95276.diff 33 Files Affected: - (modified) clang/include/clang/AST/ASTContext.h (+2) - (modified) clang/include/clang/AST/Type.h (+3) - (modified) clang/include/clang/AST/TypeProperties.td (+4) - (added) clang/include/clang/Basic/AMDGPUTypes.def (+21) - (modified) clang/include/clang/Basic/Builtins.def (+1) - (modified) clang/include/clang/Basic/BuiltinsAMDGPU.def (+2) - (modified) clang/include/clang/Serialization/ASTBitCodes.h (+4-1) - (modified) clang/lib/AST/ASTContext.cpp (+15) - (modified) clang/lib/AST/ASTImporter.cpp (+4) - (modified) clang/lib/AST/ExprConstant.cpp (+2) - (modified) clang/lib/AST/ItaniumMangle.cpp (+6) - (modified) clang/lib/AST/MicrosoftMangle.cpp (+2) - (modified) clang/lib/AST/NSAPI.cpp (+2) - (modified) clang/lib/AST/PrintfFormatString.cpp (+2) - (modified) clang/lib/AST/Type.cpp (+9) - (modified) clang/lib/AST/TypeLoc.cpp (+2) - (modified) clang/lib/CodeGen/CGBuiltin.cpp (+9) - (modified) clang/lib/CodeGen/CGDebugInfo.cpp (+9-1) - (modified) clang/lib/CodeGen/CGDebugInfo.h (+2) - (modified) clang/lib/CodeGen/CodeGenTypes.cpp (+2) - (modified) clang/lib/CodeGen/ItaniumCXXABI.cpp (+2) - (modified) clang/lib/Index/USRGeneration.cpp (+4) - (modified) clang/lib/Sema/Sema.cpp (+6) - (modified) clang/lib/Sema/SemaExpr.cpp (+4) - (modified) clang/lib/Serialization/ASTCommon.cpp (+5) - (modified) clang/lib/Serialization/ASTReader.cpp (+5) - (added) clang/test/AST/ast-dump-amdgpu-types.c (+10) - (added) clang/test/CodeGen/amdgpu-buffer-rsrc-type-debug-info.c (+9) - (added) clang/test/CodeGenCXX/amdgpu-buffer-rsrc-typeinfo.cpp (+9) - (added) clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl (+26) - (added) clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl (+14) - (added) clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp (+16) - (modified) clang/tools/libclang/CIndex.cpp (+2) ``````````diff diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index a1d1d1c51cd41..2328141b27e79 100644 --- a/clang/include/clang/AST/ASTContext.h +++ b/clang/include/clang/AST/ASTContext.h @@ -1147,6 +1147,8 @@ class ASTContext : public RefCountedBase<ASTContext> { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) CanQualType SingletonId; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) CanQualType SingletonId; +#include "clang/Basic/AMDGPUTypes.def" // Types for deductions in C++0x [stmt.ranged]'s desugaring. Built on demand. mutable QualType AutoDeductTy; // Deduction against 'auto'. diff --git a/clang/include/clang/AST/Type.h b/clang/include/clang/AST/Type.h index 9eb3f6c09e3d3..cbcd6d0f97efe 100644 --- a/clang/include/clang/AST/Type.h +++ b/clang/include/clang/AST/Type.h @@ -3015,6 +3015,9 @@ class BuiltinType : public Type { // WebAssembly reference types #define WASM_TYPE(Name, Id, SingletonId) Id, #include "clang/Basic/WebAssemblyReferenceTypes.def" +// AMDGPU types +#define AMDGPU_TYPE(Name, Id, SingletonId) Id, +#include "clang/Basic/AMDGPUTypes.def" // All other builtin types #define BUILTIN_TYPE(Id, SingletonId) Id, #define LAST_BUILTIN_TYPE(Id) LastKind = Id diff --git a/clang/include/clang/AST/TypeProperties.td b/clang/include/clang/AST/TypeProperties.td index 40dd16f080e2e..aba14b222a03a 100644 --- a/clang/include/clang/AST/TypeProperties.td +++ b/clang/include/clang/AST/TypeProperties.td @@ -861,6 +861,10 @@ let Class = BuiltinType in { case BuiltinType::ID: return ctx.SINGLETON_ID; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(NAME, ID, SINGLETON_ID) \ + case BuiltinType::ID: return ctx.SINGLETON_ID; +#include "clang/Basic/AMDGPUTypes.def" + #define BUILTIN_TYPE(ID, SINGLETON_ID) \ case BuiltinType::ID: return ctx.SINGLETON_ID; #include "clang/AST/BuiltinTypes.def" diff --git a/clang/include/clang/Basic/AMDGPUTypes.def b/clang/include/clang/Basic/AMDGPUTypes.def new file mode 100644 index 0000000000000..226e75480037c --- /dev/null +++ b/clang/include/clang/Basic/AMDGPUTypes.def @@ -0,0 +1,21 @@ +//===-- AMDGPUTypes.def - Metadata about AMDGPU types -----------*- C++ -*-===// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// +// This file defines various AMDGPU builtin types. +// +//===----------------------------------------------------------------------===// + +#ifndef AMDGPU_OPAQUE_TYPE +#define AMDGPU_OPAQUE_TYPE(Name, MangledName, Id, SingletonId) \ + AMDGPU_TYPE(Name, Id, SingletonId) +#endif + +AMDGPU_OPAQUE_TYPE("__buffer_rsrc_t", "__buffer_rsrc_t", AMDGPUBufferRsrc, AMDGPUBufferRsrcTy) + +#undef AMDGPU_TYPE +#undef AMDGPU_OPAQUE_TYPE diff --git a/clang/include/clang/Basic/Builtins.def b/clang/include/clang/Basic/Builtins.def index f356f881d5ef9..c4dc0627f2a0f 100644 --- a/clang/include/clang/Basic/Builtins.def +++ b/clang/include/clang/Basic/Builtins.def @@ -33,6 +33,7 @@ // q -> Scalable vector, followed by the number of elements and the base type. // Q -> target builtin type, followed by a character to distinguish the builtin type // Qa -> AArch64 svcount_t builtin type. +// Qb -> AMDGPU __buffer_rsrc_t builtin type. // E -> ext_vector, followed by the number of elements and the base type. // X -> _Complex, followed by the base type. // Y -> ptrdiff_t diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 9e6800ea814a0..a73e63355cfd7 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -148,6 +148,8 @@ BUILTIN(__builtin_amdgcn_qsad_pk_u16_u8, "WUiWUiUiWUi", "nc") BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi", "nc") BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc") +BUILTIN(__builtin_amdgcn_make_buffer_rsrc, "Qbv*sii", "nc") + //===----------------------------------------------------------------------===// // Ballot builtins. //===----------------------------------------------------------------------===// diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h index 69b71e409dbe2..986548a51cbd4 100644 --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1091,6 +1091,9 @@ enum PredefinedTypeIDs { // \brief WebAssembly reference types with auto numeration #define WASM_TYPE(Name, Id, SingletonId) PREDEF_TYPE_##Id##_ID, #include "clang/Basic/WebAssemblyReferenceTypes.def" +// \brief AMDGPU types with auto numeration +#define AMDGPU_TYPE(Name, Id, SingletonId) PREDEF_TYPE_##Id##_ID, +#include "clang/Basic/AMDGPUTypes.def" /// The placeholder type for unresolved templates. PREDEF_TYPE_UNRESOLVED_TEMPLATE, @@ -1103,7 +1106,7 @@ enum PredefinedTypeIDs { /// /// Type IDs for non-predefined types will start at /// NUM_PREDEF_TYPE_IDs. -const unsigned NUM_PREDEF_TYPE_IDS = 503; +const unsigned NUM_PREDEF_TYPE_IDS = 504; // Ensure we do not overrun the predefined types we reserved // in the enum PredefinedTypeIDs above. diff --git a/clang/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp index bf74e56a14799..a49e01d5535cd 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1384,6 +1384,12 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, #include "clang/Basic/WebAssemblyReferenceTypes.def" } + if (Target.getTriple().isAMDGPU()) { +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + InitBuiltinType(SingletonId, BuiltinType::Id); +#include "clang/Basic/AMDGPUTypes.def" + } + // Builtin type for __objc_yes and __objc_no ObjCBuiltinBoolTy = (Target.useSignedCharForObjCBool() ? SignedCharTy : BoolTy); @@ -2200,6 +2206,9 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T) const { Align = 8; \ break; #include "clang/Basic/WebAssemblyReferenceTypes.def" + case BuiltinType::AMDGPUBufferRsrc: + Width = 0; + Align = 128; } break; case Type::ObjCObjectPointer: @@ -8155,6 +8164,8 @@ static char getObjCEncodingForPrimitiveType(const ASTContext *C, #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" { DiagnosticsEngine &Diags = C->getDiagnostics(); unsigned DiagID = Diags.getCustomDiagID(DiagnosticsEngine::Error, @@ -11516,6 +11527,10 @@ static QualType DecodeTypeFromStr(const char *&Str, const ASTContext &Context, Type = Context.SveCountTy; break; } + case 'b': { + Type = Context.AMDGPUBufferRsrcTy; + break; + } default: llvm_unreachable("Unexpected target builtin type"); } diff --git a/clang/lib/AST/ASTImporter.cpp b/clang/lib/AST/ASTImporter.cpp index 02cd4ed9a6cac..1b67feaae8874 100644 --- a/clang/lib/AST/ASTImporter.cpp +++ b/clang/lib/AST/ASTImporter.cpp @@ -1099,6 +1099,10 @@ ExpectedType ASTNodeImporter::VisitBuiltinType(const BuiltinType *T) { case BuiltinType::Id: \ return Importer.getToContext().SingletonId; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + case BuiltinType::Id: \ + return Importer.getToContext().SingletonId; +#include "clang/Basic/AMDGPUTypes.def" #define SHARED_SINGLETON_TYPE(Expansion) #define BUILTIN_TYPE(Id, SingletonId) \ case BuiltinType::Id: return Importer.getToContext().SingletonId; diff --git a/clang/lib/AST/ExprConstant.cpp b/clang/lib/AST/ExprConstant.cpp index af1f18aa8ef24..818e16292a5ef 100644 --- a/clang/lib/AST/ExprConstant.cpp +++ b/clang/lib/AST/ExprConstant.cpp @@ -11859,6 +11859,8 @@ GCCTypeClass EvaluateBuiltinClassifyType(QualType T, #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" return GCCTypeClass::None; case BuiltinType::Dependent: diff --git a/clang/lib/AST/ItaniumMangle.cpp b/clang/lib/AST/ItaniumMangle.cpp index eac1801445255..217fed7da37e5 100644 --- a/clang/lib/AST/ItaniumMangle.cpp +++ b/clang/lib/AST/ItaniumMangle.cpp @@ -3423,6 +3423,12 @@ void CXXNameMangler::mangleType(const BuiltinType *T) { Out << 'u' << type_name.size() << type_name; \ break; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_OPAQUE_TYPE(InternalName, MangledName, Id, SingletonId) \ + case BuiltinType::Id: \ + type_name = MangledName; \ + Out << 'u' << type_name.size() << type_name; \ + break; +#include "clang/Basic/AMDGPUTypes.def" } } diff --git a/clang/lib/AST/MicrosoftMangle.cpp b/clang/lib/AST/MicrosoftMangle.cpp index ffc5d2d4cd8fc..a92682b3bd36b 100644 --- a/clang/lib/AST/MicrosoftMangle.cpp +++ b/clang/lib/AST/MicrosoftMangle.cpp @@ -2611,6 +2611,8 @@ void MicrosoftCXXNameMangler::mangleType(const BuiltinType *T, Qualifiers, #include "clang/Basic/PPCTypes.def" #define RVV_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/RISCVVTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::ShortAccum: case BuiltinType::Accum: case BuiltinType::LongAccum: diff --git a/clang/lib/AST/NSAPI.cpp b/clang/lib/AST/NSAPI.cpp index 2d16237f5325a..48d1763125e6c 100644 --- a/clang/lib/AST/NSAPI.cpp +++ b/clang/lib/AST/NSAPI.cpp @@ -453,6 +453,8 @@ NSAPI::getNSNumberFactoryMethodKind(QualType T) const { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::BoundMember: case BuiltinType::UnresolvedTemplate: case BuiltinType::Dependent: diff --git a/clang/lib/AST/PrintfFormatString.cpp b/clang/lib/AST/PrintfFormatString.cpp index dd3b38fabb550..3031d76abbd75 100644 --- a/clang/lib/AST/PrintfFormatString.cpp +++ b/clang/lib/AST/PrintfFormatString.cpp @@ -865,6 +865,8 @@ bool PrintfSpecifier::fixType(QualType QT, const LangOptions &LangOpt, #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" #define SIGNED_TYPE(Id, SingletonId) #define UNSIGNED_TYPE(Id, SingletonId) #define FLOATING_TYPE(Id, SingletonId) diff --git a/clang/lib/AST/Type.cpp b/clang/lib/AST/Type.cpp index 33acae2cbafac..bc1c692ee9997 100644 --- a/clang/lib/AST/Type.cpp +++ b/clang/lib/AST/Type.cpp @@ -2446,6 +2446,9 @@ bool Type::isSizelessBuiltinType() const { // WebAssembly reference types #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_OPAQUE_TYPE(Name, MangledName, Id, SingletonId) \ + case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" return true; default: return false; @@ -3509,6 +3512,10 @@ StringRef BuiltinType::getName(const PrintingPolicy &Policy) const { case Id: \ return Name; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + case Id: \ + return Name; +#include "clang/Basic/AMDGPUTypes.def" } llvm_unreachable("Invalid builtin type."); @@ -4778,6 +4785,8 @@ bool Type::canHaveNullability(bool ResultIfUnknown) const { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::BuiltinFn: case BuiltinType::NullPtr: case BuiltinType::IncompleteMatrixIdx: diff --git a/clang/lib/AST/TypeLoc.cpp b/clang/lib/AST/TypeLoc.cpp index 9dd90d9bf4e54..33e6ccbadc12d 100644 --- a/clang/lib/AST/TypeLoc.cpp +++ b/clang/lib/AST/TypeLoc.cpp @@ -428,6 +428,8 @@ TypeSpecifierType BuiltinTypeLoc::getWrittenTypeSpec() const { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::BuiltinFn: case BuiltinType::IncompleteMatrixIdx: case BuiltinType::ArraySection: diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp index 511e1fd4016d7..dcc6b2a912b13 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -19082,6 +19082,15 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType}); return Builder.CreateCall(F, {Arg}); } + case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: { + llvm::Value *Base = EmitScalarExpr(E->getArg(0)); + llvm::Value *Stride = EmitScalarExpr(E->getArg(1)); + llvm::Value *Num = EmitScalarExpr(E->getArg(2)); + llvm::Value *Flags = EmitScalarExpr(E->getArg(3)); + Function *F = + CGM.getIntrinsic(Intrinsic::amdgcn_make_buffer_rsrc, {Base->getType()}); + return Builder.CreateCall(F, {Base, Stride, Num, Flags}); + } default: return nullptr; } diff --git a/clang/lib/CodeGen/CGDebugInfo.cpp b/clang/lib/CodeGen/CGDebugInfo.cpp index 99e12da0081af..fca83609e904b 100644 --- a/clang/lib/CodeGen/CGDebugInfo.cpp +++ b/clang/lib/CodeGen/CGDebugInfo.cpp @@ -865,7 +865,15 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType *BT) { return SingletonId; \ } #include "clang/Basic/WebAssemblyReferenceTypes.def" - +#define AMDGPU_OPAQUE_TYPE(Name, MangledName, Id, SingletonId) \ + case BuiltinType::Id: { \ + if (!SingletonId) \ + SingletonId = \ + DBuilder.createForwardDecl(llvm::dwarf::DW_TAG_structure_type, \ + MangledName, TheCU, TheCU->getFile(), 0); \ + return SingletonId; \ + } +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::UChar: case BuiltinType::Char_U: Encoding = llvm::dwarf::DW_ATE_unsigned_char; diff --git a/clang/lib/CodeGen/CGDebugInfo.h b/clang/lib/CodeGen/CGDebugInfo.h index 8fe738be21568..26aa0fd9b2000 100644 --- a/clang/lib/CodeGen/CGDebugInfo.h +++ b/clang/lib/CodeGen/CGDebugInfo.h @@ -83,6 +83,8 @@ class CGDebugInfo { #include "clang/Basic/OpenCLExtensionTypes.def" #define WASM_TYPE(Name, Id, SingletonId) llvm::DIType *SingletonId = nullptr; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) llvm::DIType *SingletonId = nullptr; +#include "clang/Basic/AMDGPUTypes.def" /// Cache of previously constructed Types. llvm::DenseMap<const void *, llvm::TrackingMDRef> TypeCache; diff --git a/clang/lib/CodeGen/CodeGenTypes.cpp b/clang/lib/CodeGen/CodeGenTypes.cpp index 0a926e4ac27fe..d157eb00aa593 100644 --- a/clang/lib/CodeGen/CodeGenTypes.cpp +++ b/clang/lib/CodeGen/CodeGenTypes.cpp @@ -533,6 +533,8 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) { llvm_unreachable("Unexpected wasm reference builtin type!"); \ } break; #include "clang/Basic/WebAssemblyReferenceTypes.def" + case BuiltinType::AMDGPUBufferRsrc: + return llvm::PointerType::get(getLLVMContext(), /*AddressSpace=*/8); case BuiltinType::Dependent: #define BUILTIN_TYPE(Id, SingletonId) #define PLACEHOLDER_TYPE(Id, SingletonId) \ diff --git a/clang/lib/CodeGen/ItaniumCXXABI.cpp b/clang/lib/CodeGen/ItaniumCXXABI.cpp index 8427286dee887..da3bdeb900756 100644 --- a/clang/lib/CodeGen/ItaniumCXXABI.cpp +++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp @@ -3360,6 +3360,8 @@ static bool TypeInfoIsInStandardLibrary(const BuiltinType *Ty) { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::ShortAccum: case BuiltinType::Accum: case BuiltinType::LongAccum: diff --git a/clang/lib/Index/USRGeneration.cpp b/clang/lib/Index/USRGeneration.cpp index 31c4a3345c09d..0ec1cc0d4154c 100644 --- a/clang/lib/Index/USRGeneration.cpp +++ b/clang/lib/Index/USRGeneration.cpp @@ -770,6 +770,10 @@ void USRGenerator::VisitType(QualType T) { case BuiltinType::Id: \ Out << "@BT@" << Name; break; #include "clang/Basic/RISCVVTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + case BuiltinType::Id: \ + Out << "@BT@" << #Name; break; +#include "clang/Basic/AMDGPUTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" case BuiltinType::ShortAccum: diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp index a612dcd4b4d03..add48578a18fd 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -497,6 +497,12 @@ void Sema::Initialize() { #include "clang/Basic/WebAssemblyReferenceTypes.def" } + if (Context.getTargetInfo().getTriple().isAMDGPU()) { +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + addImplicitTypedef(Name, Context.SingletonId); +#include "clang/Basic/AMDGPUTypes.def" + } + if (Context.getTargetInfo().hasBuiltinMSVaList()) { DeclarationName MSVaList = &Context.Idents.get("__builtin_ms_va_list"); if (IdResolver.begin(MSVaList) == IdResolver.end()) diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp index 44f886bf54e3a..76b9a54e2b1b4 100644 --- a/clang/lib/Sema/SemaExpr.cpp +++ b/clang/lib/Sema/SemaExpr.cpp @@ -6169,6 +6169,8 @@ static bool isPlaceholderToRemoveAsArg(QualType type) { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" #define PLACEHOLDER_TYPE(ID, SINGLETON_ID) #define BUILTIN_TYPE(ID, SINGLETON_ID) case BuiltinType::ID: #include "clang/AST/BuiltinTypes.def" @@ -20994,6 +20996,8 @@ ExprResult Sema::CheckPlaceholderExpr(Expr *E) { #include "clang/Basic/RISCVVTypes.def" #define WASM_TYPE(Name, Id, SingletonId) case BuiltinType::Id: #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) case BuiltinType::Id: +#include "clang/Basic/AMDGPUTypes.def" #define BUILTIN_TYPE(Id, SingletonId) case BuiltinType::Id: #define PLACEHOLDER_TYPE(Id, SingletonId) #include "clang/AST/BuiltinTypes.def" diff --git a/clang/lib/Serialization/ASTCommon.cpp b/clang/lib/Serialization/ASTCommon.cpp index bc662a87a7bf3..3385cb8aad7e4 100644 --- a/clang... [truncated] `````````` </details> https://github.com/llvm/llvm-project/pull/95276 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
