diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
index b3211db8df2dcd..94b44714b530de 100644
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6545,11 +6545,9 @@ ExprResult Sema::BuildCallExpr(Scope *Scope, Expr *Fn, SourceLocation LParenLoc,
         auto ArgPtTy = ArgTy->getPointeeType();
         auto ArgAS = ArgPtTy.getAddressSpace();
 
-        // Add address space cast if target address spaces are different
-        if ((ArgAS != LangAS::Default &&
-             getASTContext().getTargetAddressSpace(ArgAS) !=
-                 getASTContext().getTargetAddressSpace(ParamAS)) ||
-            ParamAS == LangAS::Default)
+        // Only allow implicit casting from a non-default address space pointee
+        // type to a default address space pointee type
+        if (ArgAS != LangAS::Default || ParamAS == LangAS::Default)
           continue;
 
         // First, ensure that the Arg is an RValue.
diff --git a/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu b/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
deleted file mode 100644
index d15953b3cacaa0..00000000000000
--- a/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
+++ /dev/null
@@ -1,20 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \
-// RUN:  -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device -emit-llvm %s \
-// RUN:  -o - | FileCheck %s
-
-#define __device__ __attribute__((device))
-typedef __attribute__((address_space(3))) float *LP;
-
-// CHECK-LABEL: test_ds_atomic_add_f32
-// CHECK: %[[ADDR_ADDR:.*]] = alloca float*, align 8, addrspace(5)
-// CHECK: %[[ADDR_ADDR_ASCAST_PTR:.*]] = addrspacecast float* addrspace(5)* %[[ADDR_ADDR]] to float**
-// CHECK: store float* %addr, float** %[[ADDR_ADDR_ASCAST_PTR]], align 8
-// CHECK: %[[ADDR_ADDR_ASCAST:.*]] = load float*, float** %[[ADDR_ADDR_ASCAST_PTR]], align 8
-// CHECK: %[[AS_CAST:.*]] = addrspacecast float* %[[ADDR_ADDR_ASCAST]] to float addrspace(3)*
-// CHECK: %3 = call contract float @llvm.amdgcn.ds.fadd.f32(float addrspace(3)* %[[AS_CAST]]
-// CHECK: %4 = load float*, float** %rtn.ascast, align 8
-// CHECK: store float %3, float* %4, align 4
-__device__ void test_ds_atomic_add_f32(float *addr, float val) {
-  float *rtn;
-  *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);
-}
diff --git a/clang/test/SemaCUDA/builtins-unsafe-atomics-gfx90a.cu b/clang/test/SemaCUDA/builtins-unsafe-atomics-gfx90a.cu
deleted file mode 100644
index 6f1484c68ec711..00000000000000
--- a/clang/test/SemaCUDA/builtins-unsafe-atomics-gfx90a.cu
+++ /dev/null
@@ -1,12 +0,0 @@
-// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -target-cpu gfx90a -x hip \
-// RUN:  -aux-triple x86_64-unknown-linux-gnu -fcuda-is-device %s \
-// RUN:  -fsyntax-only -verify
-// expected-no-diagnostics
-
-#define __device__ __attribute__((device))
-typedef __attribute__((address_space(3))) float *LP;
-
-__device__ void test_ds_atomic_add_f32(float *addr, float val) {
-  float *rtn;
-  *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);
-}