[hip] Enable pointer argument lowering through coercing type.

[hliao]

It happens that Sam has a similar patch of this one. After discussion, we agreed that this patch addresses more cases found in the workloads. Thank Sam for the test case.

[yaxunl]

we need a test for byval struct and array. better to have a struct containing an array which contain another struct which contains a pointer. thanks.

[tra]

clang/lib/CodeGen/TargetInfo.cpp

7688

Nit: for lower -> for lowering or that lowers

7690

I don't think you need a class here -- it just complicates calling of coerce().
I'd just make coerce() a member function.

7696

Nit: VM in VMCtx is not useful. Ctx or LLVMCtx would be better, IMO.

[arsenm]

clang/lib/CodeGen/CGCall.cpp

1308–1310

I would somewhat prefer 2 dyn_cast and getAddressSpace, this is essentially isa + cast combo

clang/lib/CodeGen/TargetInfo.cpp

7719

No tests with arrays or structs?

It's also not immediately obvious to me that this optimization is still valid if the pointer is buried in a struct

[hliao]

add the test case for struct.

[hliao]

clang/lib/CodeGen/TargetInfo.cpp

7719

the original generic kernel pointer promotion to a global one only handles the pointer directly passed. From a critical workload, I found quite a few cases where the global pointers are passed through a by-val struct. We didn't handle that yet. With this case, we could start to handle that.

7719

struct tests are added. From test cases, it seems to me that arry is not passed by value. I need to double-confirm.

[hliao]

revise code following reviwers' comments.

[tra]

clang/lib/CodeGen/TargetInfo.cpp

7689

Now it could use a more descriptive name, too. :-)

You can now also make DefaultAS/GlobalAS into local variables as you have access to getContext() here.

[hliao]

• revise member function name.
• add the test case for by-val array types.

[hliao]

clang/lib/CodeGen/TargetInfo.cpp

7689

name is changed but I want to leave DefaultAS and GlobalAS as parameters as they may vary from HIP to OpenCL and different targets. Even though it may be rare case, I want to avoid careless errors.

7719

a test case for arrary types is added.

[tra]

clang/lib/CodeGen/TargetInfo.cpp

7689

You may not need it, ever and it would be easy to add, but I'll leave it up to you.

If you do want to keep them as parameters you may want to consider renaming them to FromAS/ToAS.
There's nothing in the code that has anything to do with whether they are for generic/specific address space and the function name does not indicate the direction of coercion between them. It's very easy to pass them in the wrong order and not notice it. Making them local variables would avoid it. Giving names some sort of 'directionality' would at least give user a hint what goes where, even if it would not prevent making the error.

[hliao]

revise parameter names

[hliao]

clang/lib/CodeGen/TargetInfo.cpp

7689

From the target device side, we have generic and global addresses. But, at the language level, we have opencl_global and cuda_device. Even though they map into the same address space, it would be very confusing if they are misused to initialize that address space numbers. That's why the original helper makes more sense to me and makes the code more readable. Anyway, I change the parameter names to give a clear direction.

[tra]

Thank you!

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

2–3

Perhaps we should add host-side test, too to make sure the pointers there do remain generic.

[hliao]

Add host-side checks.

[yaxunl]

I am a little bit concerned that user may have such code:

struct A { int *p; }
__global__ kernel(A a) {
  int x;
  a.p = &x;
  f(a);
}

@arsenm what happens if a private pointer is mis-used as a global pointer?

I am wondering if we should coerce byval struct kernel arg to global only if they are const, e.g.

__global__ kernel(const A a);

I understand this may lose performance. Or should we introduce an option to let user disable coerce of non-const struct kernel arg to global.

[hliao]

In D69826#1734296, @yaxunl wrote:

I am a little bit concerned that user may have such code:

struct A { int *p; }
__global__ kernel(A a) {
  int x;
  a.p = &x;
  f(a);
}

@arsenm what happens if a private pointer is mis-used as a global pointer?

I am wondering if we should coerce byval struct kernel arg to global only if they are const, e.g.

__global__ kernel(const A a);

I understand this may lose performance. Or should we introduce an option to let user disable coerce of non-const struct kernel arg to global.

This should not be a concern. The coercing is only applied to the parameter itself. Within the function body, we still use the original struct A. The preparation in function prolog will copy that coerced argument into the original one (alloca-ed.) The modification of that parameter later will be applied to the original one due to the by-val nature.

A modified version of your code is compiled into the following code at O0:

define protected amdgpu_kernel void @_Z3foo1A(%struct.A.coerce %a.coerce) #0 {
entry:
  %a = alloca %struct.A, align 8, addrspace(5)
  %a1 = addrspacecast %struct.A addrspace(5)* %a to %struct.A*
  %x = alloca i32, align 4, addrspace(5)
  %x.ascast = addrspacecast i32 addrspace(5)* %x to i32*
  %agg.tmp = alloca %struct.A, align 8, addrspace(5)
  %agg.tmp.ascast = addrspacecast %struct.A addrspace(5)* %agg.tmp to %struct.A*
  %0 = bitcast %struct.A* %a1 to %struct.A.coerce*
  %1 = getelementptr inbounds %struct.A.coerce, %struct.A.coerce* %0, i32 0, i32 0
  %2 = extractvalue %struct.A.coerce %a.coerce, 0
  store i32 addrspace(1)* %2, i32 addrspace(1)** %1, align 8
  %3 = getelementptr inbounds %struct.A.coerce, %struct.A.coerce* %0, i32 0, i32 1
  %4 = extractvalue %struct.A.coerce %a.coerce, 1
  store i32 addrspace(1)* %4, i32 addrspace(1)** %3, align 8
  %p = getelementptr inbounds %struct.A, %struct.A* %a1, i32 0, i32 0
  store i32* %x.ascast, i32** %p, align 8
  %5 = bitcast %struct.A* %agg.tmp.ascast to i8*
  %6 = bitcast %struct.A* %a1 to i8*
  call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %5, i8* align 8 %6, i64 16, i1 false)
  %7 = getelementptr inbounds %struct.A, %struct.A* %agg.tmp.ascast, i32 0, i32 0
  %8 = load i32*, i32** %7, align 8
  %9 = getelementptr inbounds %struct.A, %struct.A* %agg.tmp.ascast, i32 0, i32 1
  %10 = load i32*, i32** %9, align 8
  call void @_Z1f1A(i32* %8, i32* %10) #3
  ret void
}

The modification of parameter a is applied the alloca-ed one.

[yaxunl]

In D69826#1734324, @hliao wrote:

In D69826#1734296, @yaxunl wrote:

I am a little bit concerned that user may have such code:

struct A { int *p; }
__global__ kernel(A a) {
  int x;
  a.p = &x;
  f(a);
}

@arsenm what happens if a private pointer is mis-used as a global pointer?

I am wondering if we should coerce byval struct kernel arg to global only if they are const, e.g.

__global__ kernel(const A a);

I understand this may lose performance. Or should we introduce an option to let user disable coerce of non-const struct kernel arg to global.

This should not be a concern. The coercing is only applied to the parameter itself. Within the function body, we still use the original struct A. The preparation in function prolog will copy that coerced argument into the original one (alloca-ed.) The modification of that parameter later will be applied to the original one due to the by-val nature.

A modified version of your code is compiled into the following code at O0:

define protected amdgpu_kernel void @_Z3foo1A(%struct.A.coerce %a.coerce) #0 {
entry:
  %a = alloca %struct.A, align 8, addrspace(5)
  %a1 = addrspacecast %struct.A addrspace(5)* %a to %struct.A*
  %x = alloca i32, align 4, addrspace(5)
  %x.ascast = addrspacecast i32 addrspace(5)* %x to i32*
  %agg.tmp = alloca %struct.A, align 8, addrspace(5)
  %agg.tmp.ascast = addrspacecast %struct.A addrspace(5)* %agg.tmp to %struct.A*
  %0 = bitcast %struct.A* %a1 to %struct.A.coerce*
  %1 = getelementptr inbounds %struct.A.coerce, %struct.A.coerce* %0, i32 0, i32 0
  %2 = extractvalue %struct.A.coerce %a.coerce, 0
  store i32 addrspace(1)* %2, i32 addrspace(1)** %1, align 8
  %3 = getelementptr inbounds %struct.A.coerce, %struct.A.coerce* %0, i32 0, i32 1
  %4 = extractvalue %struct.A.coerce %a.coerce, 1
  store i32 addrspace(1)* %4, i32 addrspace(1)** %3, align 8
  %p = getelementptr inbounds %struct.A, %struct.A* %a1, i32 0, i32 0
  store i32* %x.ascast, i32** %p, align 8
  %5 = bitcast %struct.A* %agg.tmp.ascast to i8*
  %6 = bitcast %struct.A* %a1 to i8*
  call void @llvm.memcpy.p0i8.p0i8.i64(i8* align 8 %5, i8* align 8 %6, i64 16, i1 false)
  %7 = getelementptr inbounds %struct.A, %struct.A* %agg.tmp.ascast, i32 0, i32 0
  %8 = load i32*, i32** %7, align 8
  %9 = getelementptr inbounds %struct.A, %struct.A* %agg.tmp.ascast, i32 0, i32 1
  %10 = load i32*, i32** %9, align 8
  call void @_Z1f1A(i32* %8, i32* %10) #3
  ret void
}

The modification of parameter a is applied the alloca-ed one.

OK. Thanks for clarification.

clang/lib/CodeGen/CGCall.cpp

static void CreateCoercedStore(llvm::Value *Src,
 
   uint64_t SrcSize = CGF.CGM.getDataLayout().getTypeAllocSize(SrcTy);
 
   if (llvm::StructType *DstSTy = dyn_cast<llvm::StructType>(DstTy)) {
     Dst = EnterStructPointerForCoercedAccess(Dst, DstSTy, SrcSize, CGF);
     DstTy = Dst.getType()->getElementType();
   }
 
+  if (isa<llvm::PointerType>(SrcTy) &&
+      isa<llvm::PointerType>(DstTy) &&
+      SrcTy->getPointerAddressSpace() != DstTy->getPointerAddressSpace()) {
+    Src = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(Src, DstTy);
+    CGF.Builder.CreateStore(Src, Dst, DstIsVolatile);
+    return;
+  }
+
   // If the source and destination are integer or pointer types, just do an
   // extension or truncation to the desired type.
   if ((isa<llvm::IntegerType>(SrcTy) || isa<llvm::PointerType>(SrcTy)) &&
       (isa<llvm::IntegerType>(DstTy) || isa<llvm::PointerType>(DstTy))) {
     Src = CoerceIntOrPtrToIntOrPtr(Src, DstTy, CGF);
     CGF.Builder.CreateStore(Src, Dst, DstIsVolatile);
     return;
   }

clang/lib/CodeGen/TargetInfo.cpp

private:
   static const unsigned MaxNumRegsForArgsRet = 16;
 
   unsigned numRegsForType(QualType Ty) const;
 
   bool isHomogeneousAggregateBaseType(QualType Ty) const override;
   bool isHomogeneousAggregateSmallEnough(const Type *Base,
                                          uint64_t Members) const override;
 
+  // Coercion type builder for lower HIP pointer argument from generic pointer
+  // to global pointer.
+  class CoerceGenericPointerTypeBuilder {
+    llvm::LLVMContext &Context;
+    unsigned DefaultAS;
+    unsigned GlobalAS;
+
+  public:
+    CoerceGenericPointerTypeBuilder(llvm::LLVMContext &VMCtx, unsigned DAS,
+                                    unsigned GAS)
+        : Context(VMCtx), DefaultAS(DAS), GlobalAS(GAS) {}
+
+    llvm::Type *coerce(llvm::Type *Ty) {
+      // Structure types.
+      if (auto STy = dyn_cast<llvm::StructType>(Ty)) {
+        SmallVector<llvm::Type *, 8> EltTys;
+        bool Changed = false;
+        for (auto T : STy->elements()) {
+          auto NT = coerce(T);
+          EltTys.push_back(NT);
+          Changed |= (NT != T);
+        }
+        // Skip if there is no change in element types.
+        if (!Changed)
+          return STy;
+        if (STy->hasName())
+          return llvm::StructType::create(
+              EltTys, (STy->getName() + ".coerce").str(), STy->isPacked());
+        return llvm::StructType::get(Context, EltTys, STy->isPacked());
+      }
+      // Arrary types.
+      if (auto ATy = dyn_cast<llvm::ArrayType>(Ty)) {
+        auto T = ATy->getElementType();
+        auto NT = coerce(T);
+        // Skip if there is no change in that element type.
+        if (NT == T)
+          return ATy;
+        return llvm::ArrayType::get(NT, ATy->getNumElements());
+      }
+      // Single value types.
+      if (Ty->isPointerTy() && Ty->getPointerAddressSpace() == DefaultAS)
+        return llvm::PointerType::get(
+            cast<llvm::PointerType>(Ty)->getElementType(), GlobalAS);
+      return Ty;
+    }
+  };
+
 public:
   explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) :
     DefaultABIInfo(CGT) {}
 
   ABIArgInfo classifyReturnType(QualType RetTy) const;
   ABIArgInfo classifyKernelArgumentType(QualType Ty) const;
   ABIArgInfo classifyArgumentType(QualType Ty, unsigned &NumRegsLeft) const;
 

 
 /// For kernels all parameters are really passed in a special buffer. It doesn't
 /// make sense to pass anything byval, so everything must be direct.
 ABIArgInfo AMDGPUABIInfo::classifyKernelArgumentType(QualType Ty) const {
   Ty = useFirstFieldIfTransparentUnion(Ty);
 
   // TODO: Can we omit empty structs?
 
-  // Coerce single element structs to its element.
+  llvm::Type *LTy = nullptr;
   if (const Type *SeltTy = isSingleElementStruct(Ty, getContext()))
-    return ABIArgInfo::getDirect(CGT.ConvertType(QualType(SeltTy, 0)));
+    LTy = CGT.ConvertType(QualType(SeltTy, 0));
+
+  if (getContext().getLangOpts().HIP) {
+    if (!LTy)
+      LTy = CGT.ConvertType(Ty);
+    CoerceGenericPointerTypeBuilder Builder(getVMContext(),
+        getContext().getTargetAddressSpace(LangAS::Default),
+        getContext().getTargetAddressSpace(LangAS::cuda_device));
+    LTy = Builder.coerce(LTy);
+  }
 
   // If we set CanBeFlattened to true, CodeGen will expand the struct to its
   // individual elements, which confuses the Clover OpenCL backend; therefore we
   // have to set it to false here. Other args of getDirect() are just defaults.
-  return ABIArgInfo::getDirect(nullptr, 0, nullptr, false);
+  return ABIArgInfo::getDirect(LTy, 0, nullptr, false);
 }
 
 ABIArgInfo AMDGPUABIInfo::classifyArgumentType(QualType Ty,
                                                unsigned &NumRegsLeft) const {
   assert(NumRegsLeft <= MaxNumRegsForArgsRet && "register estimate underflow");
 
   Ty = useFirstFieldIfTransparentUnion(Ty);
 

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

+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device \
+// RUN:     -emit-llvm -x hip %s -o - | FileCheck %s
+#include "Inputs/cuda.h"
+// CHECK: define amdgpu_kernel void  @_Z7kernel1Pi(i32 addrspace(1)* %x.coerce)
+__global__ void kernel1(int *x) {
+  x[0]++;
+}
+
+// CHECK: define amdgpu_kernel void  @_Z7kernel2Ri(i32 addrspace(1)* dereferenceable(4) %x.coerce)
+__global__ void kernel2(int &x) {
+  x++;
+}
+
+// CHECK: define amdgpu_kernel void  @_Z7kernel3PU3AS2iPU3AS1i(i32 addrspace(2)* %x, i32 addrspace(1)* %y)
+__global__ void kernel3(__attribute__((address_space(2))) int *x,
+                        __attribute__((address_space(1))) int *y) {
+  y[0] = x[0];
+}
+
+// CHECK: define void @_Z4funcPi(i32* %x)
+__device__ void func(int *x) {
+  x[0]++;
+}