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

Depends on #94830.

Author: shiltian

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

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.
 //===----------------------------------------------------------------------===//

clang/lib/AST/ASTContext.cpp

@@ -11527,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");
     }

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;
   }

clang/test/CodeGenOpenCL/builtins-amdgcn-make-buffer-rsrc.cl

@@ -0,0 +1,14 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu verde -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu tonga -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-unknown-unknown -target-cpu gfx1100 -emit-llvm -o - %s | FileCheck %s
+
+// CHECK-LABEL: @test_amdgcn_make_buffer_rsrc(
+// CHECK-NEXT:  entry:
+// CHECK-NEXT:    [[TMP0:%.*]] = tail call ptr addrspace(8) @llvm.amdgcn.make.buffer.rsrc.p5(ptr addrspace(5) [[P:%.*]], i16 4, i32 [[NUM:%.*]], i32 [[FLAGS:%.*]])
+// 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);
+ }