https://github.com/shiltian updated
https://github.com/llvm/llvm-project/pull/95276
>From 013a40d474e3acaa7a090d5e279f2d8a2f18fbd8 Mon Sep 17 00:00:00 2001
From: Shilei Tian <i...@tianshilei.me>
Date: Mon, 17 Jun 2024 18:48:33 -0400
Subject: [PATCH 1/2] [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 | 16 ++++
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 | 6 ++
clang/lib/AST/TypeLoc.cpp | 2 +
clang/lib/CodeGen/CGDebugInfo.cpp | 11 ++-
clang/lib/CodeGen/CGDebugInfo.h | 2 +
clang/lib/CodeGen/CodeGenTypes.cpp | 5 ++
clang/lib/CodeGen/ItaniumCXXABI.cpp | 2 +
clang/lib/Index/USRGeneration.cpp | 5 ++
clang/lib/Sema/Sema.cpp | 8 ++
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 | 8 ++
.../amdgpu-buffer-rsrc-typeinfo.cpp | 9 ++
.../CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl | 82 +++++++++++++++++++
clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp | 17 ++++
clang/test/SemaHIP/amdgpu-buffer-rsrc.hip | 20 +++++
clang/test/SemaOpenCL/amdgpu-buffer-rsrc.cl | 12 +++
clang/test/SemaOpenMP/amdgpu-buffer-rsrc.cpp | 17 ++++
clang/tools/libclang/CIndex.cpp | 2 +
32 files changed, 299 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/SemaCXX/amdgpu-buffer-rsrc.cpp
create mode 100644 clang/test/SemaHIP/amdgpu-buffer-rsrc.hip
create mode 100644 clang/test/SemaOpenCL/amdgpu-buffer-rsrc.cl
create mode 100644 clang/test/SemaOpenMP/amdgpu-buffer-rsrc.cpp
diff --git a/clang/include/clang/AST/ASTContext.h
b/clang/include/clang/AST/ASTContext.h
index 53ece996769a8..4d1f440506e09 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 fab233b62d8d1..61246479188e9 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..e0d7be470a325
--- /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_PTR_TYPE
+#define AMDGPU_OPAQUE_PTR_TYPE(Name, MangledName, AS, Width, Align, Id,
SingletonId) \
+ AMDGPU_TYPE(Name, Id, SingletonId)
+#endif
+
+AMDGPU_OPAQUE_PTR_TYPE("__amdgpu_buffer_rsrc_t", "__amdgpu_buffer_rsrc_t", 8,
128, 128, AMDGPUBufferRsrc, AMDGPUBufferRsrcTy)
+
+#undef AMDGPU_TYPE
+#undef AMDGPU_OPAQUE_PTR_TYPE
diff --git a/clang/include/clang/Serialization/ASTBitCodes.h
b/clang/include/clang/Serialization/ASTBitCodes.h
index a4728b1c06b3f..24e616f76b9af 100644
--- a/clang/include/clang/Serialization/ASTBitCodes.h
+++ b/clang/include/clang/Serialization/ASTBitCodes.h
@@ -1097,6 +1097,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,
@@ -1109,7 +1112,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 34aa399fda2f8..d389ef12468ee 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -1384,6 +1384,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);
@@ -2200,6 +2207,13 @@ TypeInfo ASTContext::getTypeInfoImpl(const Type *T)
const {
Align = 8;
\
break;
#include "clang/Basic/WebAssemblyReferenceTypes.def"
+#define AMDGPU_OPAQUE_PTR_TYPE(NAME, MANGLEDNAME, AS, WIDTH, ALIGN, ID,
\
+ SINGLETONID)
\
+ case BuiltinType::ID:
\
+ Width = WIDTH;
\
+ Align = ALIGN;
\
+ break;
+#include "clang/Basic/AMDGPUTypes.def"
}
break;
case Type::ObjCObjectPointer:
@@ -8168,6 +8182,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 3a6c8b4f82cca..dd355a7125c5a 100644
--- a/clang/lib/AST/ExprConstant.cpp
+++ b/clang/lib/AST/ExprConstant.cpp
@@ -11814,6 +11814,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 ed9e6eeb36c75..203db72c43733 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_TYPE(Name, Id, SingletonId)
\
+ case BuiltinType::Id:
\
+ type_name = Name;
\
+ 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 a863ec7a529b9..d87be5f2043a9 100644
--- a/clang/lib/AST/MicrosoftMangle.cpp
+++ b/clang/lib/AST/MicrosoftMangle.cpp
@@ -2612,6 +2612,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..656b733a13b0e 100644
--- a/clang/lib/AST/Type.cpp
+++ b/clang/lib/AST/Type.cpp
@@ -3509,6 +3509,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 +4782,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 572ba84d22ef5..a072475ba7705 100644
--- a/clang/lib/CodeGen/CGDebugInfo.cpp
+++ b/clang/lib/CodeGen/CGDebugInfo.cpp
@@ -865,7 +865,16 @@ llvm::DIType *CGDebugInfo::CreateType(const BuiltinType
*BT) {
return SingletonId;
\
}
#include "clang/Basic/WebAssemblyReferenceTypes.def"
-
+#define AMDGPU_OPAQUE_PTR_TYPE(Name, MangledName, AS, Width, Align, 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 272c8d6e75965..2731c627d9dc3 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..d823c336e39bf 100644
--- a/clang/lib/CodeGen/CodeGenTypes.cpp
+++ b/clang/lib/CodeGen/CodeGenTypes.cpp
@@ -533,6 +533,11 @@ llvm::Type *CodeGenTypes::ConvertType(QualType T) {
llvm_unreachable("Unexpected wasm reference builtin type!");
\
} break;
#include "clang/Basic/WebAssemblyReferenceTypes.def"
+#define AMDGPU_OPAQUE_PTR_TYPE(Name, MangledName, AS, Width, Align, Id,
\
+ SingletonId)
\
+ case BuiltinType::Id:
\
+ return llvm::PointerType::get(getLLVMContext(), AS);
+#include "clang/Basic/AMDGPUTypes.def"
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 5a3e83de625c9..01a735c1437e1 100644
--- a/clang/lib/CodeGen/ItaniumCXXABI.cpp
+++ b/clang/lib/CodeGen/ItaniumCXXABI.cpp
@@ -3362,6 +3362,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..5036ddee35fd1 100644
--- a/clang/lib/Index/USRGeneration.cpp
+++ b/clang/lib/Index/USRGeneration.cpp
@@ -772,6 +772,11 @@ void USRGenerator::VisitType(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:
\
+ Out << "@BT@" << #Name;
\
+ break;
+#include "clang/Basic/AMDGPUTypes.def"
case BuiltinType::ShortAccum:
Out << "@BT@ShortAccum"; break;
case BuiltinType::Accum:
diff --git a/clang/lib/Sema/Sema.cpp b/clang/lib/Sema/Sema.cpp
index 907a05a5d1b49..069978c1b4023 100644
--- a/clang/lib/Sema/Sema.cpp
+++ b/clang/lib/Sema/Sema.cpp
@@ -497,6 +497,14 @@ void Sema::Initialize() {
#include "clang/Basic/WebAssemblyReferenceTypes.def"
}
+ if (Context.getTargetInfo().getTriple().isAMDGPU() ||
+ (Context.getAuxTargetInfo() &&
+ 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 4db8b4130c3c7..a31cefc540cbf 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"
@@ -21004,6 +21006,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 a2c322087fd1e..0810d720bb4e0 100644
--- a/clang/lib/Serialization/ASTReader.cpp
+++ b/clang/lib/Serialization/ASTReader.cpp
@@ -7401,6 +7401,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..e032d678f1a09
--- /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
__amdgpu_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 __amdgpu_buffer_rsrc_t /dev/null | sed -e "s/ <undeserialized
declarations>//" -e "s/ imported//" | FileCheck %s
+
+// CHECK: TypedefDecl {{.*}} implicit __amdgpu_buffer_rsrc_t
+// CHECK-NEXT: -BuiltinType {{.*}} '__amdgpu_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..c266fa83e4b62
--- /dev/null
+++ b/clang/test/CodeGen/amdgpu-buffer-rsrc-type-debug-info.c
@@ -0,0 +1,8 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn -emit-llvm -o - %s -debug-info-kind=limited
2>&1 | FileCheck %s
+
+// CHECK: name: "__amdgpu_buffer_rsrc_t",{{.*}}baseType: ![[BT:[0-9]+]]
+// CHECK: [[BT]] = !DICompositeType(tag: DW_TAG_structure_type, name:
"__amdgpu_buffer_rsrc_t", {{.*}} flags: DIFlagFwdDecl)
+void test_locals(void) {
+ __amdgpu_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..a44e7dc5efe6a
--- /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(__amdgpu_buffer_rsrc_t);
+
+// CHECK-DAG: @_ZTSu22__amdgpu_buffer_rsrc_t = {{.*}}
c"u22__amdgpu_buffer_rsrc_t\00"
+// CHECK-DAG: @_ZTIu22__amdgpu_buffer_rsrc_t = {{.*}}
@_ZTVN10__cxxabiv123__fundamental_type_infoE, {{.*}}
@_ZTSu22__amdgpu_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..69dabda08fba6
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/amdgcn-buffer-rsrc-type.cl
@@ -0,0 +1,82 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --function-signature
+ // REQUIRES: amdgpu-registered-target
+ // RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde
-emit-llvm -o - %s | FileCheck %s
+
+typedef struct AA_ty {
+ int x;
+ __amdgpu_buffer_rsrc_t r;
+} AA;
+
+AA getAA(void *p);
+__amdgpu_buffer_rsrc_t getBufferImpl(void *p);
+void consumeBuffer(__amdgpu_buffer_rsrc_t);
+
+// CHECK-LABEL: define {{[^@]+}}@getBuffer
+// CHECK-SAME: (ptr addrspace(5) noundef [[P:%.*]]) local_unnamed_addr
#[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[CALL:%.*]] = tail call ptr addrspace(8) @getBufferImpl(ptr
addrspace(5) noundef [[P]]) #[[ATTR2:[0-9]+]]
+// CHECK-NEXT: ret ptr addrspace(8) [[CALL]]
+//
+__amdgpu_buffer_rsrc_t getBuffer(void *p) {
+ return getBufferImpl(p);
+}
+
+// CHECK-LABEL: define {{[^@]+}}@consumeBufferPtr
+// CHECK-SAME: (ptr addrspace(5) noundef readonly [[P:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TOBOOL_NOT:%.*]] = icmp eq ptr addrspace(5) [[P]],
addrspacecast (ptr null to ptr addrspace(5))
+// CHECK-NEXT: br i1 [[TOBOOL_NOT]], label [[IF_END:%.*]], label
[[IF_THEN:%.*]]
+// CHECK: if.then:
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr addrspace(8), ptr addrspace(5)
[[P]], align 16, !tbaa [[TBAA4:![0-9]+]]
+// CHECK-NEXT: tail call void @consumeBuffer(ptr addrspace(8) [[TMP0]])
#[[ATTR2]]
+// CHECK-NEXT: br label [[IF_END]]
+// CHECK: if.end:
+// CHECK-NEXT: ret void
+//
+void consumeBufferPtr(__amdgpu_buffer_rsrc_t *p) {
+ if (p)
+ consumeBuffer(*p);
+}
+
+// CHECK-LABEL: define {{[^@]+}}@test
+// CHECK-SAME: (ptr addrspace(5) noundef readonly [[A:%.*]])
local_unnamed_addr #[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = load i32, ptr addrspace(5) [[A]], align 16,
!tbaa [[TBAA8:![0-9]+]]
+// CHECK-NEXT: [[TOBOOL_NOT:%.*]] = icmp eq i32 [[TMP0]], 0
+// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq ptr addrspace(5) [[A]],
addrspacecast (ptr null to ptr addrspace(5))
+// CHECK-NEXT: [[OR_COND:%.*]] = or i1 [[TOBOOL_NOT_I]], [[TOBOOL_NOT]]
+// CHECK-NEXT: br i1 [[OR_COND]], label [[IF_END:%.*]], label
[[IF_THEN_I:%.*]]
+// CHECK: if.then.i:
+// CHECK-NEXT: [[R:%.*]] = getelementptr inbounds i8, ptr addrspace(5)
[[A]], i32 16
+// CHECK-NEXT: [[TMP1:%.*]] = load ptr addrspace(8), ptr addrspace(5)
[[R]], align 16, !tbaa [[TBAA4]]
+// CHECK-NEXT: tail call void @consumeBuffer(ptr addrspace(8) [[TMP1]])
#[[ATTR2]]
+// CHECK-NEXT: br label [[IF_END]]
+// CHECK: if.end:
+// CHECK-NEXT: ret void
+//
+void test(AA *a) {
+ if (a->x)
+ consumeBufferPtr(&(a->r));
+}
+
+// CHECK-LABEL: define {{[^@]+}}@bar
+// CHECK-SAME: (ptr addrspace(5) noundef [[P:%.*]]) local_unnamed_addr
#[[ATTR0]] {
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[CALL:%.*]] = tail call [[STRUCT_AA_TY:%.*]]
@[[GETAA:[a-zA-Z0-9_$\"\\.-]*[a-zA-Z_$\"\\.-][a-zA-Z0-9_$\"\\.-]*]](ptr
addrspace(5) noundef [[P]]) #[[ATTR2]]
+// CHECK-NEXT: [[TMP0:%.*]] = extractvalue [[STRUCT_AA_TY]] [[CALL]], 0
+// CHECK-NEXT: [[CALL_I:%.*]] = tail call ptr addrspace(8)
@getBufferImpl(ptr addrspace(5) noundef [[P]]) #[[ATTR2]]
+// CHECK-NEXT: [[TOBOOL_NOT_I:%.*]] = icmp eq i32 [[TMP0]], 0
+// CHECK-NEXT: br i1 [[TOBOOL_NOT_I]], label [[TEST_EXIT:%.*]], label
[[IF_THEN_I_I:%.*]]
+// CHECK: if.then.i.i:
+// CHECK-NEXT: tail call void @consumeBuffer(ptr addrspace(8) [[CALL_I]])
#[[ATTR2]]
+// CHECK-NEXT: br label [[TEST_EXIT]]
+// CHECK: test.exit:
+// CHECK-NEXT: [[DOTFCA_1_INSERT:%.*]] = insertvalue [[STRUCT_AA_TY]]
[[CALL]], ptr addrspace(8) [[CALL_I]], 1
+// CHECK-NEXT: ret [[STRUCT_AA_TY]] [[DOTFCA_1_INSERT]]
+//
+AA bar(void *p) {
+ AA a = getAA(p);
+ a.r = getBuffer(p);
+ test(&a);
+ return a;
+}
diff --git a/clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp
b/clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp
new file mode 100644
index 0000000000000..80c4c519c4e6b
--- /dev/null
+++ b/clang/test/SemaCXX/amdgpu-buffer-rsrc.cpp
@@ -0,0 +1,17 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -std=gnu++11 -triple amdgcn
-Wno-unused-value %s
+
+void foo() {
+ int n = 100;
+ __amdgpu_buffer_rsrc_t v = 0; // expected-error {{cannot initialize a
variable of type '__amdgpu_buffer_rsrc_t' with an rvalue of type 'int'}}
+ static_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error {{static_cast from
'int' to '__amdgpu_buffer_rsrc_t' is not allowed}}
+ dynamic_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error {{invalid target
type '__amdgpu_buffer_rsrc_t' for dynamic_cast; target type must be a reference
or pointer type to a defined class}}
+ reinterpret_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error
{{reinterpret_cast from 'int' to '__amdgpu_buffer_rsrc_t' is not allowed}}
+ int c(v); // expected-error {{cannot initialize a variable of type 'int'
with an lvalue of type '__amdgpu_buffer_rsrc_t'}}
+ __amdgpu_buffer_rsrc_t k;
+ int *ip = (int *)k; // expected-error {{cannot cast from type
'__amdgpu_buffer_rsrc_t' to pointer type 'int *'}}
+ void *vp = (void *)k; // expected-error {{cannot cast from type
'__amdgpu_buffer_rsrc_t' to pointer type 'void *'}}
+}
+
+static_assert(sizeof(__amdgpu_buffer_rsrc_t) == 16, "wrong size");
+static_assert(alignof(__amdgpu_buffer_rsrc_t) == 16, "wrong aignment");
diff --git a/clang/test/SemaHIP/amdgpu-buffer-rsrc.hip
b/clang/test/SemaHIP/amdgpu-buffer-rsrc.hip
new file mode 100644
index 0000000000000..3e5b22dc8963d
--- /dev/null
+++ b/clang/test/SemaHIP/amdgpu-buffer-rsrc.hip
@@ -0,0 +1,20 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -fsyntax-only -verify -triple amdgcn -Wno-unused-value %s
+// RUN: %clang_cc1 -fsyntax-only -verify -triple x86_64 -aux-triple amdgcn
-Wno-unused-value %s
+
+#define __device__ __attribute__((device))
+
+__device__ void foo() {
+ int n = 100;
+ __amdgpu_buffer_rsrc_t v = 0; // expected-error {{cannot initialize a
variable of type '__amdgpu_buffer_rsrc_t' with an rvalue of type 'int'}}
+ static_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error {{static_cast from
'int' to '__amdgpu_buffer_rsrc_t' is not allowed}}
+ dynamic_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error {{invalid target
type '__amdgpu_buffer_rsrc_t' for dynamic_cast; target type must be a reference
or pointer type to a defined class}}
+ reinterpret_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error
{{reinterpret_cast from 'int' to '__amdgpu_buffer_rsrc_t' is not allowed}}
+ int c(v); // expected-error {{cannot initialize a variable of type 'int'
with an lvalue of type '__amdgpu_buffer_rsrc_t'}}
+ __amdgpu_buffer_rsrc_t k;
+ int *ip = (int *)k; // expected-error {{cannot cast from type
'__amdgpu_buffer_rsrc_t' to pointer type 'int *'}}
+ void *vp = (void *)k; // expected-error {{cannot cast from type
'__amdgpu_buffer_rsrc_t' to pointer type 'void *'}}
+}
+
+static_assert(sizeof(__amdgpu_buffer_rsrc_t) == 16, "wrong size");
+static_assert(alignof(__amdgpu_buffer_rsrc_t) == 16, "wrong aignment");
diff --git a/clang/test/SemaOpenCL/amdgpu-buffer-rsrc.cl
b/clang/test/SemaOpenCL/amdgpu-buffer-rsrc.cl
new file mode 100644
index 0000000000000..2d74835699c6d
--- /dev/null
+++ b/clang/test/SemaOpenCL/amdgpu-buffer-rsrc.cl
@@ -0,0 +1,12 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -cl-std=CL1.2 -triple amdgcn-amd-amdhsa
-Wno-unused-value %s
+// RUN: %clang_cc1 -verify -cl-std=CL2.0 -triple amdgcn-amd-amdhsa
-Wno-unused-value %s
+
+void foo() {
+ int n = 100;
+ __amdgpu_buffer_rsrc_t v = 0; // expected-error {{initializing '__private
__amdgpu_buffer_rsrc_t' with an expression of incompatible type 'int'}}
+ int c = v; // expected-error {{initializing '__private int' with an
expression of incompatible type '__private __amdgpu_buffer_rsrc_t'}}
+ __amdgpu_buffer_rsrc_t k;
+ int *ip = (int *)k; // expected-error {{operand of type
'__amdgpu_buffer_rsrc_t' where arithmetic or pointer type is required}}
+ void *vp = (void *)k; // expected-error {{operand of type
'__amdgpu_buffer_rsrc_t' where arithmetic or pointer type is required}}
+ }
diff --git a/clang/test/SemaOpenMP/amdgpu-buffer-rsrc.cpp
b/clang/test/SemaOpenMP/amdgpu-buffer-rsrc.cpp
new file mode 100644
index 0000000000000..eb6ded229a75c
--- /dev/null
+++ b/clang/test/SemaOpenMP/amdgpu-buffer-rsrc.cpp
@@ -0,0 +1,17 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=amdgcn-amd-amdhsa -triple
amdgcn-amd-amdhsa -fopenmp-is-target-device -Wno-unused-value %s
+
+void foo() {
+#pragma omp target
+ {
+ int n = 100;
+ __amdgpu_buffer_rsrc_t v = 0; // expected-error {{cannot initialize a
variable of type '__amdgpu_buffer_rsrc_t' with an rvalue of type 'int'}}
+ static_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error {{static_cast
from 'int' to '__amdgpu_buffer_rsrc_t' is not allowed}}
+ dynamic_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error {{invalid
target type '__amdgpu_buffer_rsrc_t' for dynamic_cast; target type must be a
reference or pointer type to a defined class}}
+ reinterpret_cast<__amdgpu_buffer_rsrc_t>(n); // expected-error
{{reinterpret_cast from 'int' to '__amdgpu_buffer_rsrc_t' is not allowed}}
+ int c(v); // expected-error {{cannot initialize a variable of type 'int'
with an lvalue of type '__amdgpu_buffer_rsrc_t'}}
+ __amdgpu_buffer_rsrc_t k;
+ int *ip = (int *)k; // expected-error {{cannot cast from type
'__amdgpu_buffer_rsrc_t' to pointer type 'int *'}}
+ void *vp = (void *)k; // expected-error {{cannot cast from type
'__amdgpu_buffer_rsrc_t' to pointer type 'void *'}}
+ }
+ }
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:
>From 74f1ca15746865f2f895b7dd6e20978de4e8df3c Mon Sep 17 00:00:00 2001
From: Shilei Tian <i...@tianshilei.me>
Date: Tue, 18 Jun 2024 12:16:11 -0400
Subject: [PATCH 2/2] [Clang][AMDGPU] Add a builtin for
`llvm.amdgcn.make.buffer.rsrc` intrinsic
Depends on #94830.
---
clang/include/clang/Basic/Builtins.def | 1 +
clang/include/clang/Basic/BuiltinsAMDGPU.def | 2 +
clang/lib/AST/ASTContext.cpp | 4 +
clang/lib/CodeGen/CGBuiltin.cpp | 13 +++
.../CodeGenHIP/builtins-make-buffer-rsrc.hip | 105 ++++++++++++++++++
.../builtins-amdgcn-make-buffer-rsrc.cl | 93 ++++++++++++++++
6 files changed, 218 insertions(+)
create mode 100644 clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip
create mode 100644 clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
diff --git a/clang/include/clang/Basic/Builtins.def
b/clang/include/clang/Basic/Builtins.def
index f356f881d5ef9..d2d500c990b99 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 __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/lib/AST/ASTContext.cpp b/clang/lib/AST/ASTContext.cpp
index d389ef12468ee..f7cb87ffb0326 100644
--- a/clang/lib/AST/ASTContext.cpp
+++ b/clang/lib/AST/ASTContext.cpp
@@ -11545,6 +11545,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/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index 511e1fd4016d7..ff36a03ccec84 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -615,6 +615,17 @@ static Value *emitTernaryBuiltin(CodeGenFunction &CGF,
return CGF.Builder.CreateCall(F, { Src0, Src1, Src2 });
}
+static Value *emitQuaternaryBuiltin(CodeGenFunction &CGF, const CallExpr *E,
+ unsigned IntrinsicID) {
+ llvm::Value *Src0 = CGF.EmitScalarExpr(E->getArg(0));
+ llvm::Value *Src1 = CGF.EmitScalarExpr(E->getArg(1));
+ llvm::Value *Src2 = CGF.EmitScalarExpr(E->getArg(2));
+ llvm::Value *Src3 = CGF.EmitScalarExpr(E->getArg(3));
+
+ Function *F = CGF.CGM.getIntrinsic(IntrinsicID, Src0->getType());
+ return CGF.Builder.CreateCall(F, {Src0, Src1, Src2, Src3});
+}
+
// Emit an intrinsic that has 1 float or double operand, and 1 integer.
static Value *emitFPIntBuiltin(CodeGenFunction &CGF,
const CallExpr *E,
@@ -19082,6 +19093,8 @@ 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:
+ return emitQuaternaryBuiltin(*this, E, Intrinsic::amdgcn_make_buffer_rsrc);
default:
return nullptr;
}
diff --git a/clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip
b/clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip
new file mode 100644
index 0000000000000..c1a30633f3d0a
--- /dev/null
+++ b/clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip
@@ -0,0 +1,105 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
UTC_ARGS: --version 5
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu verde -emit-llvm
-disable-llvm-optzns -fcuda-is-device -o - %s | FileCheck %s
+
+#define __device__ __attribute__((device))
+
+// CHECK-LABEL: define dso_local ptr addrspace(8)
@_Z31test_amdgcn_make_buffer_rsrc_p0Pvsii(
+// CHECK-SAME: ptr noundef [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32
noundef [[NUM:%.*]], i32 noundef [[FLAGS:%.*]]) #[[ATTR0:[0-9]+]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16,
addrspace(5)
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[STRIDE_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
+// CHECK-NEXT: [[NUM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[FLAGS_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[P_ADDR]] to ptr
+// CHECK-NEXT: [[STRIDE_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[STRIDE_ADDR]] to ptr
+// CHECK-NEXT: [[NUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[NUM_ADDR]] to ptr
+// CHECK-NEXT: [[FLAGS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[FLAGS_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i16 [[STRIDE]], ptr [[STRIDE_ADDR_ASCAST]], align 2
+// CHECK-NEXT: store i32 [[NUM]], ptr [[NUM_ADDR_ASCAST]], align 4
+// CHECK-NEXT: store i32 [[FLAGS]], ptr [[FLAGS_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[STRIDE_ADDR_ASCAST]], align 2
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[NUM_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = load i32, ptr [[FLAGS_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP4:%.*]] = call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 [[TMP1]], i32 [[TMP2]], i32
[[TMP3]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP4]]
+//
+__device__ __amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p,
short stride, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
+}
+
+// CHECK-LABEL: define dso_local ptr addrspace(8)
@_Z47test_amdgcn_make_buffer_rsrc_p0_stride_constantPvii(
+// CHECK-SAME: ptr noundef [[P:%.*]], i32 noundef [[NUM:%.*]], i32 noundef
[[FLAGS:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16,
addrspace(5)
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[NUM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[FLAGS_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[P_ADDR]] to ptr
+// CHECK-NEXT: [[NUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[NUM_ADDR]] to ptr
+// CHECK-NEXT: [[FLAGS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[FLAGS_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i32 [[NUM]], ptr [[NUM_ADDR_ASCAST]], align 4
+// CHECK-NEXT: store i32 [[FLAGS]], ptr [[FLAGS_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[NUM_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[FLAGS_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 4, i32 [[TMP1]], i32
[[TMP2]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP3]]
+//
+__device__ __amdgpu_buffer_rsrc_t
test_amdgcn_make_buffer_rsrc_p0_stride_constant(void *p, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
+}
+
+// CHECK-LABEL: define dso_local ptr addrspace(8)
@_Z44test_amdgcn_make_buffer_rsrc_p0_num_constantPvsi(
+// CHECK-SAME: ptr noundef [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32
noundef [[FLAGS:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16,
addrspace(5)
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[STRIDE_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
+// CHECK-NEXT: [[FLAGS_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[P_ADDR]] to ptr
+// CHECK-NEXT: [[STRIDE_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[STRIDE_ADDR]] to ptr
+// CHECK-NEXT: [[FLAGS_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[FLAGS_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i16 [[STRIDE]], ptr [[STRIDE_ADDR_ASCAST]], align 2
+// CHECK-NEXT: store i32 [[FLAGS]], ptr [[FLAGS_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[STRIDE_ADDR_ASCAST]], align 2
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[FLAGS_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 [[TMP1]], i32 1234, i32
[[TMP2]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP3]]
+//
+__device__ __amdgpu_buffer_rsrc_t
test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p, short stride, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);
+}
+
+// CHECK-LABEL: define dso_local ptr addrspace(8)
@_Z46test_amdgcn_make_buffer_rsrc_p0_flags_constantPvsi(
+// CHECK-SAME: ptr noundef [[P:%.*]], i16 noundef signext [[STRIDE:%.*]], i32
noundef [[NUM:%.*]]) #[[ATTR0]] {
+// CHECK-NEXT: [[ENTRY:.*:]]
+// CHECK-NEXT: [[RETVAL:%.*]] = alloca ptr addrspace(8), align 16,
addrspace(5)
+// CHECK-NEXT: [[P_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
+// CHECK-NEXT: [[STRIDE_ADDR:%.*]] = alloca i16, align 2, addrspace(5)
+// CHECK-NEXT: [[NUM_ADDR:%.*]] = alloca i32, align 4, addrspace(5)
+// CHECK-NEXT: [[RETVAL_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[RETVAL]] to ptr
+// CHECK-NEXT: [[P_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[P_ADDR]] to ptr
+// CHECK-NEXT: [[STRIDE_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[STRIDE_ADDR]] to ptr
+// CHECK-NEXT: [[NUM_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5)
[[NUM_ADDR]] to ptr
+// CHECK-NEXT: store ptr [[P]], ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: store i16 [[STRIDE]], ptr [[STRIDE_ADDR_ASCAST]], align 2
+// CHECK-NEXT: store i32 [[NUM]], ptr [[NUM_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[P_ADDR_ASCAST]], align 8
+// CHECK-NEXT: [[TMP1:%.*]] = load i16, ptr [[STRIDE_ADDR_ASCAST]], align 2
+// CHECK-NEXT: [[TMP2:%.*]] = load i32, ptr [[NUM_ADDR_ASCAST]], align 4
+// CHECK-NEXT: [[TMP3:%.*]] = call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[TMP0]], i16 [[TMP1]], i32 [[TMP2]], i32
5678)
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP3]]
+//
+__device__ __amdgpu_buffer_rsrc_t
test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p, short stride, int num) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
+}
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
new file mode 100644
index 0000000000000..2c7bc10fb609c
--- /dev/null
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl
@@ -0,0 +1,93 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -target-cpu
verde -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32
[[NUM:%.*]], i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0(void *p, short stride,
int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_stride_constant(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 4, i32 [[NUM:%.*]], i32
[[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_stride_constant(void
*p, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_num_constant(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32 1234,
i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_num_constant(void *p,
short stride, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p0_flags_constant(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr [[P:%.*]], i16 [[STRIDE:%.*]], i32
[[NUM:%.*]], i32 5678)
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p0_flags_constant(void *p,
short stride, int num) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16
[[STRIDE:%.*]], i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1(global void *p, short
stride, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_stride_constant(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16 4, i32
[[NUM:%.*]], i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_stride_constant(global
void *p, int num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, /*stride=*/4, num, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_num_constant(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16
[[STRIDE:%.*]], i32 1234, i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_num_constant(global
void *p, short stride, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, /*num=*/1234, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc_p1_flags_constant(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) [[P:%.*]], i16
[[STRIDE:%.*]], i32 [[NUM:%.*]], i32 5678)
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_rsrc_p1_flags_constant(global
void *p, short stride, int num) {
+ return __builtin_amdgcn_make_buffer_rsrc(p, stride, num, /*flags=*/5678);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_p0_nullptr(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p0(ptr null, i16 [[STRIDE:%.*]], i32 [[NUM:%.*]],
i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p0_nullptr(short stride, int
num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc((void *)0LL, stride, num, flags);
+}
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_p1_nullptr(
+// CHECK-NEXT: entry:
+// CHECK-NEXT: [[TMP0:%.*]] = tail call ptr addrspace(8)
@llvm.amdgcn.make.buffer.rsrc.p1(ptr addrspace(1) null, i16 [[STRIDE:%.*]], i32
[[NUM:%.*]], i32 [[FLAGS:%.*]])
+// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
+//
+__amdgpu_buffer_rsrc_t test_amdgcn_make_buffer_p1_nullptr(short stride, int
num, int flags) {
+ return __builtin_amdgcn_make_buffer_rsrc((global void *)0LL, stride, num,
flags);
+}
_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
