[hip] Add noalias on restrict qualified coerced hip pointers

[kerbowa]

Fix test formatting.

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2267

let's also check LangOpts.CudaIsDevice

2268

let's use getContext().getTargetAddressSpace(LangAS::Default) for addrspace 0 and getContext().getTargetAddressSpace(LangAS::cuda_device) for addrspace 1

2270

For struct containing pointers, we do recursive coercing

https://github.com/llvm/llvm-project/blob/master/clang/lib/CodeGen/TargetInfo.cpp#L8224

So solely compare element type will fail. We need to add a test case for struct containing pointer, and we need to have a recursive comparison in a similar way as above code.

[kerbowa]

clang/lib/CodeGen/CGCall.cpp

2270

I can add it, but are we sure it's what we want? I think OpenCL/hcc wont have the same behavior because of https://github.com/llvm/llvm-project/blob/master/clang/lib/CodeGen/CGCall.cpp#L2435.

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2270

Michael said without that there were extra flat load/stores. That's why he did that intentionally.

[kerbowa]

clang/lib/CodeGen/CGCall.cpp

2270

Okay. What I mean is that OpenCL/hcc don't add noalias recursively to struct members that are restrict qualified. Is there any reason not to do it?

[hliao]

Basically, I think that should be a generic issue. As argument coercing doesn't handle pointer previously, we need to address this pointer-specific issue. That is, if the original argument has qualifiers being able to map onto LLVM attributes, the coerced argument should be those qualifiers as well if the new coerced one is still a pointer as the original one. We may lose more useful attributes not limited to noalias, such as alignment, nonnull, and etc.

clang/lib/CodeGen/CGCall.cpp

2556–2563

I don't think we need to check pointer address and/or HIP specific. As the generic argument processing, if the original type has any qualifiers, the coerced type (if it has a single value as the original parameter) should have those qualifiers as well. Here, we not only miss restrict but also alignment, nonnull, and etc. It should be fixed as a generic case instead of a target- or language-specific one.

[yaxunl]

clang/lib/CodeGen/CGCall.cpp

2556–2563

I agree we should migrate other argument properties for promoted pointer-type kernel arg for HIP, and that should be possible since other than the address space change, the type is the same.

However, I am not sure if we can do that for coerced function arguments in general. It may not even be pointer any more.

[hliao]

clang/lib/CodeGen/CGCall.cpp

2556–2563

If the coerced type is still a pointer but diffs on the element type, address space, or others. They should share the same qualifiers for pointers.

[hliao]

Any more comments? As this should be a performance-critical issue, shall we get conclusion and make progress for the next step?

[kerbowa]

In D79213#2018820, @hliao wrote:

Any more comments? As this should be a performance-critical issue, shall we get conclusion and make progress for the next step?

We applied this current version of the patch internally for testing which is the reason for the lack of urgency, but hopefully we can get the final generic solution submitted upstream early this week. I believe I understand your comments and it makes sense to me that it should be safe for pointers whose coerced type is also a single pointer to retain the same attributes.

[rjmccall]

clang/lib/CodeGen/CGCall.cpp

2556–2563

If the parameter type is restrict-qualified and there's a single IR argument of pointer type, you should be able to apply NoAlias, regardless of coercion, in all language modes. You can just hoist that out of the "trivial case" block.

[hliao]

I hoisted parameter attributes preparation in D79395, which depends on D79394, a cleanup to make that hoist more straight-forward.

clang/lib/CodeGen/CGCall.cpp

assert((varType->isIntegerTy() || varType->isFloatingPointTy())
          && "unexpected promotion type");
 
   if (isa<llvm::IntegerType>(varType))
     return CGF.Builder.CreateTrunc(value, varType, "arg.unpromote");
 
   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,
                                          QualType ArgType, unsigned ArgNo) {
   // FIXME: __attribute__((nonnull)) can also be applied to:
   //   - references to pointers, where the pointee is known to be
   //     nonnull (apparently a Clang extension)
   //   - transparent unions containing pointers
case ABIArgInfo::Direct: {
       }
 
       Address Alloca = CreateMemTemp(Ty, getContext().getDeclAlign(Arg),
                                      Arg->getName());
 
       // 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());
       if (ArgI.isDirect() && ArgI.getCanBeFlattened() && STy &&
           STy->getNumElements() > 1) {
         uint64_t SrcSize = CGM.getDataLayout().getTypeAllocSize(STy);
         llvm::Type *DstTy = Ptr.getElementType();
         uint64_t DstSize = CGM.getDataLayout().getTypeAllocSize(DstTy);

clang/test/CodeGenCUDA/amdgpu-kernel-arg-pointer-type.cu

 };
 // `by-val` array is also coerced.
 // CHECK: define amdgpu_kernel void @_Z7kernel61T(%struct.T.coerce %t.coerce)
 // HOST: define void @_Z22__device_stub__kernel61T(float* %t.coerce0, float* %t.coerce1)
 __global__ void kernel6(struct T t) {
   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]++;
+}