diff --git a/clang/lib/CodeGen/CGCall.cpp b/clang/lib/CodeGen/CGCall.cpp
--- a/clang/lib/CodeGen/CGCall.cpp
+++ b/clang/lib/CodeGen/CGCall.cpp
@@ -2259,6 +2259,18 @@
   return CGF.Builder.CreateFPCast(value, varType, "arg.unpromote");
 }
 
+/// Returns true if the argument is a generic HIP pointer that was coerced to a
+/// global pointer.
+bool isCoercedHIPGlobalPointer(CodeGenFunction &CGF,
+                               const LangOptions &LangOpts,
+                               const ABIArgInfo &ArgI, const QualType &Ty) {
+  return LangOpts.HIP && isa<llvm::PointerType>(ArgI.getCoerceToType()) &&
+         ArgI.getCoerceToType()->getPointerAddressSpace() == 1 &&
+         CGF.ConvertType(Ty)->getPointerAddressSpace() == 0 &&
+         ArgI.getCoerceToType()->getPointerElementType() ==
+             CGF.ConvertType(Ty)->getPointerElementType();
+}
+
 /// Returns the attribute (either parameter attribute, or function
 /// attribute), which declares argument ArgNo to be non-null.
 static const NonNullAttr *getNonNullAttr(const Decl *FD, const ParmVarDecl *PVD,
@@ -2541,6 +2553,14 @@
       // Pointer to store into.
       Address Ptr = emitAddressAtOffset(*this, Alloca, ArgI);
 
+      // Restrict qualified HIP pointers that were coerced to global pointers
+      // can be marked with the noalias attribute.
+      if (isCoercedHIPGlobalPointer(*this, getLangOpts(), ArgI, Ty) &&
+          Arg->getType().isRestrictQualified()) {
+        auto AI = cast<llvm::Argument>(FnArgs[FirstIRArg]);
+        AI->addAttr(llvm::Attribute::NoAlias);
+      }
+
       // Fast-isel and the optimizer generally like scalar values better than
       // FCAs, so we flatten them if this is safe to do for this argument.
       llvm::StructType *STy = dyn_cast<llvm::StructType>(ArgI.getCoerceToType());
diff --git a/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu b/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
--- a/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
+++ b/clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
@@ -67,3 +67,10 @@
   t.x[0][0] += 1.f;
   t.x[1][0] += 2.f;
 }
+
+// Check that coerced pointers retain the noalias attribute when qualified with __restrict.
+// CHECK: define amdgpu_kernel void @_Z7kernel7Pi(i32 addrspace(1)* noalias %x.coerce)
+// HOST: define void @_Z22__device_stub__kernel7Pi(i32* noalias %x)
+__global__ void kernel7(int* __restrict x) {
+  x[0]++;
+}