[clang] Allow pointers from any address space to gfx90a builtins

[gandhi21299]

Summary

Explicitly address space cast the pointer if its address space differs
from the address space of the parameter.

[yaxunl]

clang/test/CodeGenHIP/unsafe-atomic-ops-gfx90a.hip

14

I think we want to check the pointer argument passed to the atomic intrinsics to make sure they have addr space cast if needed. That's what your patch changes.

[gandhi21299]

• corrected test, address space cast should go under unsafeAtomicAdd(...) codegen

[yaxunl]

clang/lib/CodeGen/CGBuiltin.cpp

15764

In https://reviews.llvm.org/D102405 you already inserted an addr space cast in AST when you pass a pointer to default addr space to a builtin expecting a pointer to non-default addr space. This addr space cast in AST will be emitted as addr space cast in IR. Why do you still need to insert addr space cast here?

clang/lib/CodeGen/CGBuiltin.cpp

case AMDGPU::BI__builtin_amdgcn_ds_atomic_fadd_f32: {
     }
     llvm::Value *Addr = EmitScalarExpr(E->getArg(0));
     llvm::Value *Val = EmitScalarExpr(E->getArg(1));
     llvm::Constant *ZeroI32 = llvm::ConstantInt::getIntegerValue(
         llvm::Type::getInt32Ty(getLLVMContext()), APInt(32, 0, true));
     llvm::Constant *ZeroI1 = llvm::ConstantInt::getIntegerValue(
         llvm::Type::getInt1Ty(getLLVMContext()), APInt(1, 0));
     llvm::Function *F = CGM.getIntrinsic(IID, {ArgTy});
+    auto *AddrParamTy = F->getArg(0)->getType();
+    auto *AddrTy = Addr->getType();
+    assert(AddrTy->isPointerTy() && "Argument is not of pointer type");
+    if (AddrParamTy->getPointerAddressSpace() !=
+            AddrTy->getPointerAddressSpace() &&
+        AddrParamTy->getPointerElementType() == AddrTy->getPointerElementType())
+      Addr = Builder.CreateAddrSpaceCast(Addr, AddrParamTy);
     return Builder.CreateCall(F, {Addr, Val, ZeroI32, ZeroI32, ZeroI1});
   }
   case AMDGPU::BI__builtin_amdgcn_read_exec: {
     CallInst *CI = cast<CallInst>(
       EmitSpecialRegisterBuiltin(*this, E, Int64Ty, Int64Ty, NormalRead, "exec"));
     CI->setConvergent();
     return CI;
   }

clang/test/CodeGenHIP/unsafe-atomic-ops-gfx90a.hip

+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -munsafe-fp-atomics -target-cpu gfx90a -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip -emit-llvm -fcuda-is-device -o - %s | FileCheck %s
+
+#define __global__ __attribute__((global))
+#define __device__ __attribute__((device))
+
+// CHECK-LABEL: @_Z15unsafeAtomicAddPff(float* %addr, float %value
+__device__ inline float unsafeAtomicAdd(float* addr, float value) {
+  // CHECK: %[[ADDR_ADDR:.*]] = alloca float*, align 8, addrspace(5)
+  // CHECK: %[[ADDR_ADDR_ASCAST:.*]] = addrspacecast float* addrspace(5)* %[[ADDR_ADDR]] to float**
+  // CHECK: %[[ADDR_PTR:.*]] = load float*, float** %[[ADDR_ADDR_ASCAST]], align 8
+  // CHECK: %[[ADDR:.*]] = addrspacecast float* %[[ADDR_PTR]] to float addrspace(3)*
+  // CHECK: call contract float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %[[ADDR]]
+  return __builtin_amdgcn_ds_atomic_fadd_f32(addr, value);
+}
+
+__global__ void test_global_atomic_add_f32(float *val){
+  float *rtn;
+  *rtn = unsafeAtomicAdd(val, 1.0);
+}