Skip to content

[Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic #95276

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 1 commit into from
Jun 20, 2024
Merged

[Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic #95276

merged 1 commit into from
Jun 20, 2024

Conversation

shiltian
Copy link
Contributor

@shiltian shiltian commented Jun 12, 2024 •
edited
Loading

Depends on #94830.

@llvmbot llvmbot added clang Clang issues not falling into any other category backend:AMDGPU clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:modules C++20 modules and Clang Header Modules clang:codegen IR generation bugs: mangling, exceptions, etc. debuginfo labels Jun 12, 2024
@shiltian shiltian requested review from arsenm and yxsamliu June 12, 2024 17:43
@llvmbot
Copy link
Member

llvmbot commented Jun 12, 2024 •
edited
Loading

@llvm/pr-subscribers-clang-codegen
@llvm/pr-subscribers-clang-modules
@llvm/pr-subscribers-clang

@llvm/pr-subscribers-debuginfo

Author: Shilei Tian (shiltian)

Changes
  • [Clang][AMDGPU] Add a new builtin type for buffer rsrc
  • [Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic

Patch is 28.13 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/95276.diff

33 Files Affected:

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

@shiltian shiltian changed the title buffer rsrc builtin [Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic Jun 12, 2024
@github-actions GitHub Actions
Copy link

github-actions bot commented Jun 12, 2024 •
edited
Loading

✅ With the latest revision this PR passed the C/C++ code formatter.

arsenm
arsenm reviewed Jun 12, 2024
View reviewed changes
clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl Outdated
// CHECK-NEXT: ret ptr addrspace(8) [[TMP0]]
//
__buffer_rsrc_t test_amdgcn_make_buffer_rsrc(void *p, int num, int flags) {
return __builtin_amdgcn_make_buffer_rsrc(p, 4, num, flags);
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Test more combinations of arguments. All arguments can be variable, but you have a constant here

Also test with different address space pointers. We also probably should have a HIP test. Ideally we would error for private/local inputs.

Also test null and literal pointer inputs?

@arsenm arsenm requested a review from krzysz00 June 12, 2024 17:51
@shiltian shiltian force-pushed the buffer-rsrc-builtin branch from d42a42d to e53f057 Compare June 12, 2024 18:41
@yxsamliu
Copy link
Collaborator

I am wondering whether prefix the builtin type with __amdgcn_ would be better since I envision risk of conflicting with reserved names of other compilers or standard libraries.

@shiltian
Copy link
Contributor Author

shiltian commented Jun 12, 2024 •
edited
Loading

I am wondering whether prefix the builtin type with __amdgcn_ would be better since I envision risk of conflicting with reserved names of other compilers or standard libraries.

In the patch where the type was introduced we had a brief back-and-forth. I checked the reference type WASM introduced and they don't have prefix. I don't think in the future we'd have a cross-platform/-compiler type called __buffer_rsrc_t, and if it happens, it is not supposed to have __ prefix. However, I'm by no means a language expert, so I'm fine if we really want to add that.

@yxsamliu
Copy link
Collaborator

I am wondering whether prefix the builtin type with __amdgcn_ would be better since I envision risk of conflicting with reserved names of other compilers or standard libraries.

In the patch where the type was introduced we had a brief back-and-forth. I checked the reference type WASM introduced and they don't have prefix. I don't think in the future we'd have a cross-platform/-compiler type called __buffer_rsrc_t, and if it happens, it is not supposed to have __ prefix. However, I'm by no means a language expert, so I'm fine if we really want to add that.

we are introducing __buffer_rsrc_t in global namespace, which is seen in any other namespace. Imagine some libstdc++ or libc++ header files use the same name in some namespaces and a HIP program includes these header files, there may be compilation error.

A search of __buffer shows libstdc++ and libc++ do use names starting with __buffer:

https://github.com/search?q=repo%3Agcc-mirror%2Fgcc%20path%3A%2F%5Elibstdc%5C%2B%5C%2B-v3%5C%2Finclude%5C%2F%2F%20__buffer&type=code

https://github.com/search?q=repo%3Allvm%2Fllvm-project+path%3A%2F%5Elibcxx%5C%2F%2F+__buffer&type=code&p=1

I understand the chance of conflict is low. It may be like the chance of hitting by a meteor. However, if we prefix with __amdgcn_, there is no such risk. And we have the benefit to clearly indicate it is a amdgcn target-specific type.

@arsenm
Copy link
Contributor

arsenm commented Jun 13, 2024

I understand the chance of conflict is low. It may be like the chance of hitting by a meteor. However, if we prefix with __amdgcn_, there is no such risk. And we have the benefit to clearly indicate it is a amdgcn target-specific type.

Should use amdgpu

@shiltian shiltian force-pushed the buffer-rsrc-builtin branch 2 times, most recently from df6804b to bfc92ef Compare June 18, 2024 01:51
arsenm
arsenm reviewed Jun 18, 2024
View reviewed changes
@shiltian shiltian force-pushed the buffer-rsrc-builtin branch from bfc92ef to 74f1ca1 Compare June 18, 2024 16:17
@shiltian shiltian mentioned this pull request Jun 18, 2024
Merged
@shiltian shiltian force-pushed the buffer-rsrc-builtin branch from 74f1ca1 to 5bf43d8 Compare June 19, 2024 00:49
@shiltian shiltian changed the title [Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic [Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intrinsic Jun 19, 2024
@shiltian
Copy link
Contributor Author

ping

@shiltian shiltian merged commit e3eb12c into llvm:main Jun 20, 2024
7 checks passed
@shiltian shiltian deleted the buffer-rsrc-builtin branch June 20, 2024 15:01
AlexisPerry pushed a commit to llvm-project-tlp/llvm-project that referenced this pull request Jul 9, 2024
@shiltian @AlexisPerry
[Clang][AMDGPU] Add a builtin for llvm.amdgcn.make.buffer.rsrc intr…
8a19adc 
…insic (llvm#95276)

Depends on llvm#94830.
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Reviewers

@arsenm arsenm arsenm left review comments

@yxsamliu yxsamliu yxsamliu approved these changes

@Sirraide Sirraide Awaiting requested review from Sirraide

@delcypher delcypher Awaiting requested review from delcypher

@ilya-biryukov ilya-biryukov Awaiting requested review from ilya-biryukov

@krzysz00 krzysz00 Awaiting requested review from krzysz00

Assignees
No one assigned
Labels
backend:AMDGPU clang:codegen IR generation bugs: mangling, exceptions, etc. clang:frontend Language frontend issues, e.g. anything involving "Sema" clang:modules C++20 modules and Clang Header Modules clang Clang issues not falling into any other category debuginfo
Projects
None yet
Milestone
No milestone
Development

Successfully merging this pull request may close these issues.
5 participants
@shiltian @llvmbot @yxsamliu @arsenm @krzysz00