https://github.com/shiltian updated https://github.com/llvm/llvm-project/pull/94830
>From cbb274f279c918bd0aaa601f87e50a2d10b71c25 Mon Sep 17 00:00:00 2001 From: Shilei Tian <i...@tianshilei.me> Date: Wed, 12 Jun 2024 15:46:55 -0400 Subject: [PATCH] [Clang][AMDGPU] Add a new builtin type for buffer rsrc --- clang/include/clang/AST/ASTContext.h | 2 ++ clang/include/clang/AST/Type.h | 3 +++ clang/include/clang/AST/TypeProperties.td | 4 +++ clang/include/clang/Basic/AMDGPUTypes.def | 21 +++++++++++++++ .../include/clang/Serialization/ASTBitCodes.h | 5 +++- clang/lib/AST/ASTContext.cpp | 12 +++++++++ clang/lib/AST/ASTImporter.cpp | 4 +++ clang/lib/AST/ExprConstant.cpp | 2 ++ clang/lib/AST/ItaniumMangle.cpp | 6 +++++ clang/lib/AST/MicrosoftMangle.cpp | 2 ++ clang/lib/AST/NSAPI.cpp | 2 ++ clang/lib/AST/PrintfFormatString.cpp | 2 ++ clang/lib/AST/Type.cpp | 9 +++++++ clang/lib/AST/TypeLoc.cpp | 2 ++ clang/lib/CodeGen/CGDebugInfo.cpp | 10 ++++++- clang/lib/CodeGen/CGDebugInfo.h | 2 ++ clang/lib/CodeGen/CodeGenTypes.cpp | 2 ++ clang/lib/CodeGen/ItaniumCXXABI.cpp | 2 ++ clang/lib/Index/USRGeneration.cpp | 4 +++ clang/lib/Sema/Sema.cpp | 7 +++++ clang/lib/Sema/SemaExpr.cpp | 4 +++ clang/lib/Serialization/ASTCommon.cpp | 5 ++++ clang/lib/Serialization/ASTReader.cpp | 5 ++++ clang/test/AST/ast-dump-amdgpu-types.c | 10 +++++++ .../amdgpu-buffer-rsrc-type-debug-info.c | 9 +++++++ .../amdgpu-buffer-rsrc-typeinfo.cpp | 9 +++++++ .../CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl | 26 +++++++++++++++++++ clang/test/SemaCUDA/amdgpu-buffer-rsrc-ty.cu | 6 +++++ clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp | 16 ++++++++++++ clang/tools/libclang/CIndex.cpp | 2 ++ 30 files changed, 193 insertions(+), 2 deletions(-) create mode 100644 clang/include/clang/Basic/AMDGPUTypes.def create mode 100644 clang/test/AST/ast-dump-amdgpu-types.c create mode 100644 clang/test/CodeGen/amdgpu-buffer-rsrc-type-debug-info.c create mode 100644 clang/test/CodeGenCXX/amdgpu-buffer-rsrc-typeinfo.cpp create mode 100644 clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl create mode 100644 clang/test/SemaCUDA/amdgpu-buffer-rsrc-ty.cu create mode 100644 clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp diff --git a/clang/include/clang/AST/ASTContext.h b/clang/include/clang/AST/ASTContext.h index 8bce4812f0d48..0ebe04e0886d7 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/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 cd76b8aa271da..9fc3c99cadcbf 100644 --- a/clang/lib/AST/ASTContext.cpp +++ b/clang/lib/AST/ASTContext.cpp @@ -1385,6 +1385,13 @@ void ASTContext::InitBuiltinTypes(const TargetInfo &Target, #include "clang/Basic/WebAssemblyReferenceTypes.def" } + if (Target.getTriple().isAMDGPU() || + (AuxTarget && AuxTarget->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); @@ -2201,6 +2208,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: @@ -8156,6 +8166,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, 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/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..baf311d7c45a3 100644 --- a/clang/lib/Sema/Sema.cpp +++ b/clang/lib/Sema/Sema.cpp @@ -497,6 +497,13 @@ void Sema::Initialize() { #include "clang/Basic/WebAssemblyReferenceTypes.def" } + if (Context.getTargetInfo().getTriple().isAMDGPU() || + Context.getAuxTargetInfo()->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 3869e9915075b..78a8331e2d445 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/lib/Serialization/ASTCommon.cpp +++ b/clang/lib/Serialization/ASTCommon.cpp @@ -258,6 +258,11 @@ serialization::TypeIdxFromBuiltin(const BuiltinType *BT) { ID = PREDEF_TYPE_##Id##_ID; \ break; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + case BuiltinType::Id: \ + ID = PREDEF_TYPE_##Id##_ID; \ + break; +#include "clang/Basic/AMDGPUTypes.def" case BuiltinType::BuiltinFn: ID = PREDEF_TYPE_BUILTIN_FN; break; diff --git a/clang/lib/Serialization/ASTReader.cpp b/clang/lib/Serialization/ASTReader.cpp index 59338b44db32f..f44a3d630f952 100644 --- a/clang/lib/Serialization/ASTReader.cpp +++ b/clang/lib/Serialization/ASTReader.cpp @@ -7383,6 +7383,11 @@ QualType ASTReader::GetType(TypeID ID) { T = Context.SingletonId; \ break; #include "clang/Basic/WebAssemblyReferenceTypes.def" +#define AMDGPU_TYPE(Name, Id, SingletonId) \ + case PREDEF_TYPE_##Id##_ID: \ + T = Context.SingletonId; \ + break; +#include "clang/Basic/AMDGPUTypes.def" } assert(!T.isNull() && "Unknown predefined type"); diff --git a/clang/test/AST/ast-dump-amdgpu-types.c b/clang/test/AST/ast-dump-amdgpu-types.c new file mode 100644 index 0000000000000..3347c452e6d87 --- /dev/null +++ b/clang/test/AST/ast-dump-amdgpu-types.c @@ -0,0 +1,10 @@ +// REQUIRES: amdgpu-registered-target +// Test without serialization: +// RUN: %clang_cc1 -triple amdgcn -ast-dump -ast-dump-filter __buffer_rsrc_t %s | FileCheck %s +// +// Test with serialization: +// RUN: %clang_cc1 -triple amdgcn -emit-pch -o %t %s +// RUN: %clang_cc1 -x c -triple amdgcn -include-pch %t -ast-dump-all -ast-dump-filter __buffer_rsrc_t /dev/null | sed -e "s/ <undeserialized declarations>//" -e "s/ imported//" | FileCheck %s + +// CHECK: TypedefDecl {{.*}} implicit __buffer_rsrc_t +// CHECK-NEXT: -BuiltinType {{.*}} '__buffer_rsrc_t' diff --git a/clang/test/CodeGen/amdgpu-buffer-rsrc-type-debug-info.c b/clang/test/CodeGen/amdgpu-buffer-rsrc-type-debug-info.c new file mode 100644 index 0000000000000..41a78b63e8b40 --- /dev/null +++ b/clang/test/CodeGen/amdgpu-buffer-rsrc-type-debug-info.c @@ -0,0 +1,9 @@ + +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited 2>&1 | FileCheck %s + +// CHECK: name: "__buffer_rsrc_t",{{.*}}baseType: ![[BT:[0-9]+]] +// CHECK: [[BT]] = !DICompositeType(tag: DW_TAG_structure_type, name: "__buffer_rsrc_t", {{.*}} flags: DIFlagFwdDecl) +void test_locals(void) { + __buffer_rsrc_t k; +} diff --git a/clang/test/CodeGenCXX/amdgpu-buffer-rsrc-typeinfo.cpp b/clang/test/CodeGenCXX/amdgpu-buffer-rsrc-typeinfo.cpp new file mode 100644 index 0000000000000..e733db9a8070d --- /dev/null +++ b/clang/test/CodeGenCXX/amdgpu-buffer-rsrc-typeinfo.cpp @@ -0,0 +1,9 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -triple amdgcn %s -emit-llvm -o - | FileCheck %s + +namespace std { class type_info; }; + +auto &b = typeid(__buffer_rsrc_t); + +// CHECK-DAG: @_ZTSu15__buffer_rsrc_t = {{.*}} c"u15__buffer_rsrc_t\00" +// CHECK-DAG: @_ZTIu15__buffer_rsrc_t = {{.*}} @_ZTVN10__cxxabiv123__fundamental_type_infoE, {{.*}} @_ZTSu15__buffer_rsrc_t diff --git a/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl b/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl new file mode 100644 index 0000000000000..d4f8fbeb693d6 --- /dev/null +++ b/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl @@ -0,0 +1,26 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py + // REQUIRES: amdgpu-registered-target + // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s --check-prefixes=VERDE + // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -emit-llvm -o - %s | FileCheck %s --check-prefixes=GFX8 + // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s --check-prefixes=GFX11 + +void foo(__buffer_rsrc_t); + +// VERDE-LABEL: @bar( +// VERDE-NEXT: entry: +// VERDE-NEXT: tail call void @foo(ptr addrspace(8) [[RSRC:%.*]]) #[[ATTR2:[0-9]+]] +// VERDE-NEXT: ret void +// +// GFX8-LABEL: @bar( +// GFX8-NEXT: entry: +// GFX8-NEXT: tail call void @foo(ptr addrspace(8) [[RSRC:%.*]]) #[[ATTR2:[0-9]+]] +// GFX8-NEXT: ret void +// +// GFX11-LABEL: @bar( +// GFX11-NEXT: entry: +// GFX11-NEXT: tail call void @foo(ptr addrspace(8) [[RSRC:%.*]]) #[[ATTR2:[0-9]+]] +// GFX11-NEXT: ret void +// +void bar(__buffer_rsrc_t rsrc) { + foo(rsrc); +} diff --git a/clang/test/SemaCUDA/amdgpu-buffer-rsrc-ty.cu b/clang/test/SemaCUDA/amdgpu-buffer-rsrc-ty.cu new file mode 100644 index 0000000000000..7972060ed56c7 --- /dev/null +++ b/clang/test/SemaCUDA/amdgpu-buffer-rsrc-ty.cu @@ -0,0 +1,6 @@ +// RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip -aux-triple amdgcn-amd-amdhsa %s -fsyntax-only -verify + +#define __device__ __attribute__((device)) + +__device__ __buffer_rsrc_t test_buffer_rsrc_t_device() {} // expected-warning {{non-void function does not return a value}} +__buffer_rsrc_t test_buffer_rsrc_t_host() {} // expected-error {{'__buffer_rsrc_t' can only be used in device-side function}} diff --git a/clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp b/clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp new file mode 100644 index 0000000000000..6ee7a2102821f --- /dev/null +++ b/clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp @@ -0,0 +1,16 @@ +// REQUIRES: amdgpu-registered-target +// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn -Wno-unused-value %s + +void foo() { + int n = 100; + __buffer_rsrc_t v = 0; // expected-error {{cannot initialize a variable of type '__buffer_rsrc_t' with an rvalue of type 'int'}} + static_cast<__buffer_rsrc_t>(n); // expected-error {{static_cast from 'int' to '__buffer_rsrc_t' is not allowed}} + dynamic_cast<__buffer_rsrc_t>(n); // expected-error {{invalid target type '__buffer_rsrc_t' for dynamic_cast; target type must be a reference or pointer type to a defined class}} + reinterpret_cast<__buffer_rsrc_t>(n); // expected-error {{reinterpret_cast from 'int' to '__buffer_rsrc_t' is not allowed}} + int c(v); // expected-error {{cannot initialize a variable of type 'int' with an lvalue of type '__buffer_rsrc_t'}} + sizeof(__buffer_rsrc_t); // expected-error {{invalid application of 'sizeof' to sizeless type '__buffer_rsrc_t'}} + alignof(__buffer_rsrc_t); // expected-error {{invalid application of 'alignof' to sizeless type '__buffer_rsrc_t'}} + static __buffer_rsrc_t table[1]; // expected-error {{array has sizeless element type '__buffer_rsrc_t'}} + __buffer_rsrc_t k; + int *ip = (int *)k; // expected-error {{cannot cast from type '__buffer_rsrc_t' to pointer type 'int *'}} +} diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp index dcd9555e1bfcc..35312e3d2ae70 100644 --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -1643,6 +1643,8 @@ bool CursorVisitor::VisitBuiltinTypeLoc(BuiltinTypeLoc TL) { #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) #define SIGNED_TYPE(Id, SingletonId) case BuiltinType::Id: #define UNSIGNED_TYPE(Id, SingletonId) case BuiltinType::Id: _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits