diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 7e6e0f1a06046..8e35b24f1e501 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -14,12 +14,39 @@
 
 #include "mlir/IR/Value.h"
 #include "clang/Basic/TargetBuiltins.h"
+#include "llvm/Support/AMDGPUAddrSpace.h"
 #include "llvm/Support/ErrorHandling.h"
 
 using namespace clang;
 using namespace clang::CIRGen;
 using namespace cir;
 
+// Emit the `amdgcn.dispatch.ptr` intrinsic, address-space-casting the
+// result to match \p e's return type when needed.
+// If \p e is null, returns the raw AS-4 pointer.
+static mlir::Value emitAMDGPUDispatchPtr(CIRGenFunction &cgf,
+                                         const CallExpr *e = nullptr) {
+  CIRGenBuilderTy &builder = cgf.getBuilder();
+  mlir::Location loc =
+      e ? cgf.getLoc(e->getExprLoc()) : builder.getUnknownLoc();
+  // The intrinsic always returns a pointer in the constant AS.
+  mlir::Type retTy = cir::PointerType::get(
+      cir::VoidType::get(builder.getContext()),
+      cir::TargetAddressSpaceAttr::get(builder.getContext(),
+                                       llvm::AMDGPUAS::CONSTANT_ADDRESS));
+  mlir::Value call = builder.emitIntrinsicCallOp(loc, "amdgcn.dispatch.ptr",
+                                                 retTy, mlir::ValueRange{});
+  if (!e)
+    return call;
+  // Only cast when the caller-visible AS differs from the intrinsic's AS;
+  auto expectedPtrTy =
+      mlir::cast<cir::PointerType>(cgf.convertType(e->getType()));
+  auto callPtrTy = mlir::cast<cir::PointerType>(call.getType());
+  if (expectedPtrTy.getAddrSpace() == callPtrTy.getAddrSpace())
+    return call;
+  return builder.createAddrSpaceCast(loc, call, expectedPtrTy);
+}
+
 static mlir::Value emitBinaryExpMaybeConstrainedFPBuiltin(
     CIRGenFunction &cgf, const CallExpr *e, llvm::StringRef intrinsicName,
     llvm::StringRef constrainedIntrinsicName) {
@@ -268,12 +295,8 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
                      getContext().BuiltinInfo.getName(builtinId));
     return mlir::Value{};
   }
-  case AMDGPU::BI__builtin_amdgcn_dispatch_ptr: {
-    cgm.errorNYI(expr->getSourceRange(),
-                 std::string("unimplemented AMDGPU builtin call: ") +
-                     getContext().BuiltinInfo.getName(builtinId));
-    return mlir::Value{};
-  }
+  case AMDGPU::BI__builtin_amdgcn_dispatch_ptr:
+    return emitAMDGPUDispatchPtr(*this, expr);
   case AMDGPU::BI__builtin_amdgcn_logf:
   case AMDGPU::BI__builtin_amdgcn_log_bf16: {
     cgm.errorNYI(expr->getSourceRange(),
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index b3c5ac59679a3..08ae9f52cb379 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -87,3 +87,11 @@ __device__ void test_readlane(int* out, int a, int b) {
 __device__ void test_readfirstlane(int* out, int a) {
   *out = __builtin_amdgcn_readfirstlane(a);
 }
+
+// CIR-LABEL: @_Z17test_dispatch_ptr
+// CIR: %{{.*}} = cir.call_llvm_intrinsic "amdgcn.dispatch.ptr" : () -> !cir.ptr<!void, target_address_space(4)>
+// LLVM-LABEL: @_Z17test_dispatch_ptr
+// LLVM: call{{.*}} ptr addrspace(4) @llvm.amdgcn.dispatch.ptr()
+__device__ void test_dispatch_ptr(__attribute__((address_space(4))) void ** out) {
+  *out = (__attribute__((address_space(4))) void *)__builtin_amdgcn_dispatch_ptr();
+}