diff --git a/clang/lib/Sema/SemaExpr.cpp b/clang/lib/Sema/SemaExpr.cpp
--- a/clang/lib/Sema/SemaExpr.cpp
+++ b/clang/lib/Sema/SemaExpr.cpp
@@ -6547,7 +6547,9 @@
 
         // 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)
+        if (ArgAS != LangAS::Default &&
+            getASTContext().getTargetAddressSpace(ArgAS) !=
+                getASTContext().getTargetAddressSpace(ParamAS))
           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
new file mode 100644
--- /dev/null
+++ b/clang/test/CodeGenCUDA/builtins-unsafe-atomics-gfx90a.cu
@@ -0,0 +1,12 @@
+// 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
+__device__ void test_ds_atomic_add_f32(float *addr, float val) {
+  float *rtn;
+  *rtn = __builtin_amdgcn_ds_faddf((LP)addr, val, 0, 0, 0);
+}