AMDGPU/HIP: Don't replace pointer types in kernel argument structs

[arsenm]

Summary

Currently this is counterproductive and doesn't have the desired
effect. The way the promotion is handled is by reinterpreting the
pointers in memory, which ultimately results in GVN using
ptrtoint/inttoptr. This defeats InferAddressSpaces and other
optimizations, which was the point of trying to do this replacement.

[arsenm]

I think this is obsoleted now

clang/lib/CodeGen/TargetInfo.cpp

private:
 
   bool isHomogeneousAggregateBaseType(QualType Ty) const override;
   bool isHomogeneousAggregateSmallEnough(const Type *Base,
                                          uint64_t Members) const override;
 
   // Coerce HIP pointer arguments from generic pointers to global ones.
   llvm::Type *coerceKernelArgumentType(llvm::Type *Ty, unsigned FromAS,
                                        unsigned ToAS) const {
-    // 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 = coerceKernelArgumentType(T, FromAS, ToAS);
-        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(getVMContext(), EltTys, STy->isPacked());
-    }
-    // Array types.
-    if (auto ATy = dyn_cast<llvm::ArrayType>(Ty)) {
-      auto T = ATy->getElementType();
-      auto NT = coerceKernelArgumentType(T, FromAS, ToAS);
-      // 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.
+    // TODO: This should promote generic pointers in aggregates. This used to be
+    // done, but was removed due to the way memory reinterpret resulted in the
+    // optimizer introducing ptrotoint/inttoptr. This blocks the address space
+    // inference, thereby defeating the point of doing the replacement. Single
+    // value types.
     if (Ty->isPointerTy() && Ty->getPointerAddressSpace() == FromAS)
       return llvm::PointerType::get(
           cast<llvm::PointerType>(Ty)->getElementType(), ToAS);
     return Ty;
   }
 
 public:
   explicit AMDGPUABIInfo(CodeGen::CodeGenTypes &CGT) :

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

 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -emit-llvm -x hip %s -o - | FileCheck %s
 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -emit-llvm -x hip %s -o - | FileCheck -check-prefix=HOST %s
 
 #include "Inputs/cuda.h"
 
-// Coerced struct from `struct S` without all generic pointers lowered into
-// global ones.
-// CHECK: %struct.S.coerce = type { i32 addrspace(1)*, float addrspace(1)* }
-// CHECK: %struct.T.coerce = type { [2 x float addrspace(1)*] }
+// TODO: These should coerced structs from `struct S` without all generic
+// pointers lowered into global ones when the optimizer doesn't introduce
+// ptrtoint/inttoptr
+
+// CHECK: %struct.S = type { i32*, float* }
+// CHECK: %struct.T = type { [2 x float*] }
 
 // On the host-side compilation, generic pointer won't be coerced.
 // HOST-NOT: %struct.S.coerce
 // HOST-NOT: %struct.T.coerce
 
 // CHECK: define amdgpu_kernel void  @_Z7kernel1Pi(i32 addrspace(1)* %x.coerce)
 // HOST: define void @_Z22__device_stub__kernel1Pi(i32* %x)
 __global__ void kernel1(int *x) {

 }
 
 struct S {
   int *x;
   float *y;
 };
 // `by-val` struct will be coerced into a similar struct with all generic
 // pointers lowerd into global ones.
-// CHECK: define amdgpu_kernel void @_Z7kernel41S(%struct.S.coerce %s.coerce)
+// CHECK: define amdgpu_kernel void @_Z7kernel41S(%struct.S %s.coerce)
 // HOST: define void @_Z22__device_stub__kernel41S(i32* %s.coerce0, float* %s.coerce1)
 __global__ void kernel4(struct S s) {
   s.x[0]++;
   s.y[0] += 1.f;
 }
 
 // If a pointer to struct is passed, only the pointer itself is coerced into the global one.
 // CHECK: define amdgpu_kernel void @_Z7kernel5P1S(%struct.S addrspace(1)* %s.coerce)
 // HOST: define void @_Z22__device_stub__kernel5P1S(%struct.S* %s)
 __global__ void kernel5(struct S *s) {
   s->x[0]++;
   s->y[0] += 1.f;
 }
 
 struct T {
   float *x[2];
 };
 // `by-val` array is also coerced.
-// CHECK: define amdgpu_kernel void @_Z7kernel61T(%struct.T.coerce %t.coerce)
+// CHECK: define amdgpu_kernel void @_Z7kernel61T(%struct.T %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]++;
 }
+
+struct SinglePtrEltStruct {
+  int *x;
+};
+
+// CHECK: define amdgpu_kernel void @_Z14single_ptr_elt18SinglePtrEltStruct(i32 addrspace(1)* %s.coerce)
+__global__ void single_ptr_elt(struct SinglePtrEltStruct s) {
+  s.x[0]++;
+}