[clang] [Clang][AMDGPU] Add a new builtin type for buffer rsrc (PR #94830)
Shilei Tian via cfe-commits
cfe-commits at lists.llvm.org
Wed Jun 12 12:47:08 PDT 2024
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 at 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:
More information about the cfe-commits
mailing list