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 df6804bd65ee989e897d09dee419bb435ed8872e Mon Sep 17 00:00:00 2001 From: Shilei Tian <i...@tianshilei.me> Date: Wed, 12 Jun 2024 09:18:49 -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 | 9 ++ .../CodeGenHIP/builtins-make-buffer-rsrc.hip | 107 ++++++++++++++++++ .../builtins-amdgcn-make-buffer-rsrc.cl | 95 ++++++++++++++++ 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..c4dc0627f2a0f 100644 --- a/clang/include/clang/Basic/Builtins.def +++ b/clang/include/clang/Basic/Builtins.def @@ -33,6 +33,7 @@ // q -> Scalable vector, followed by the number of elements and the base type. // Q -> target builtin type, followed by a character to distinguish the builtin type // Qa -> AArch64 svcount_t builtin type. +// Qb -> AMDGPU __buffer_rsrc_t builtin type. // E -> ext_vector, followed by the number of elements and the base type. // X -> _Complex, followed by the base type. // Y -> ptrdiff_t diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.def b/clang/include/clang/Basic/BuiltinsAMDGPU.def index 9e6800ea814a0..a73e63355cfd7 100644 --- a/clang/include/clang/Basic/BuiltinsAMDGPU.def +++ b/clang/include/clang/Basic/BuiltinsAMDGPU.def @@ -148,6 +148,8 @@ BUILTIN(__builtin_amdgcn_qsad_pk_u16_u8, "WUiWUiUiWUi", "nc") BUILTIN(__builtin_amdgcn_mqsad_pk_u16_u8, "WUiWUiUiWUi", "nc") BUILTIN(__builtin_amdgcn_mqsad_u32_u8, "V4UiWUiUiV4Ui", "nc") +BUILTIN(__builtin_amdgcn_make_buffer_rsrc, "Qbv*sii", "nc") + //===----------------------------------------------------------------------===// // Ballot builtins. //===----------------------------------------------------------------------===// diff --git a/clang/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..dcc6b2a912b13 100644 --- a/clang/lib/CodeGen/CGBuiltin.cpp +++ b/clang/lib/CodeGen/CGBuiltin.cpp @@ -19082,6 +19082,15 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID, CGM.getIntrinsic(Intrinsic::amdgcn_s_sendmsg_rtn, {ResultType}); return Builder.CreateCall(F, {Arg}); } + case AMDGPU::BI__builtin_amdgcn_make_buffer_rsrc: { + llvm::Value *Base = EmitScalarExpr(E->getArg(0)); + llvm::Value *Stride = EmitScalarExpr(E->getArg(1)); + llvm::Value *Num = EmitScalarExpr(E->getArg(2)); + llvm::Value *Flags = EmitScalarExpr(E->getArg(3)); + Function *F = + CGM.getIntrinsic(Intrinsic::amdgcn_make_buffer_rsrc, {Base->getType()}); + return Builder.CreateCall(F, {Base, Stride, Num, Flags}); + } default: return nullptr; } diff --git a/clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip b/clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip new file mode 100644 index 0000000000000..0bb8b198ca7b1 --- /dev/null +++ b/clang/test/CodeGenHIP/builtins-make-buffer-rsrc.hip @@ -0,0 +1,107 @@ +// 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 +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu tonga -emit-llvm -disable-llvm-optzns -fcuda-is-device -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx1100 -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__ __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__ __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__ __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__ __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..4eb4828f39cd3 --- /dev/null +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl @@ -0,0 +1,95 @@ +// 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 +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -target-cpu tonga -emit-llvm -o - %s | FileCheck %s +// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -cl-std=CL2.0 -target-cpu gfx1100 -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]] +// +__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]] +// +__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]] +// +__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]] +// +__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]] +// +__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]] +// +__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]] +// +__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]] +// +__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]] +// +__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]] +// +__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