[PATCH] D79213: [hip] Add noalias on restrict qualified coerced hip pointers

[Austin Kerbow via Phabricator via cfe-commits]

kerbowa updated this revision to Diff 261388.
kerbowa added a comment.

Fix test formatting.


Repository:
  rG LLVM Github Monorepo

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D79213/new/

https://reviews.llvm.org/D79213

Files:
  clang/lib/CodeGen/CGCall.cpp
  clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu


Index: clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
===================================================================
--- clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu
+++ 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]++;
+}
Index: clang/lib/CodeGen/CGCall.cpp
===================================================================
--- clang/lib/CodeGen/CGCall.cpp
+++ 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());


-------------- next part --------------
A non-text attachment was scrubbed...
Name: D79213.261388.patch
Type: text/x-patch
Size: 2471 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/cfe-commits/attachments/20200430/f44d7bc0/attachment.bin>