diff --git a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
+++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp
@@ -15980,7 +15980,7 @@
   SDValue Chain = LD->getOperand(0);
   StoreSDNode *ST = dyn_cast<StoreSDNode>(Chain.getNode());
   // TODO: Relax this restriction for unordered atomics (see D66309)
-  if (!ST || !ST->isSimple())
+  if (!ST || !ST->isSimple() || ST->getAddressSpace() != LD->getAddressSpace())
     return SDValue();
 
   EVT LDType = LD->getValueType(0);
diff --git a/llvm/test/CodeGen/NVPTX/chain-different-as.ll b/llvm/test/CodeGen/NVPTX/chain-different-as.ll
new file mode 100644
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/chain-different-as.ll
@@ -0,0 +1,10 @@
+; RUN: llc < %s -mtriple=nvptx64 | FileCheck %s
+define i64 @test() nounwind readnone {
+  %addr0 = inttoptr i64 1 to i64*
+; CHECK: mov.u64 %[[reg:rd[0-9]+]], 1
+; CHECK: ld.global.u64  %{{rd[0-9]+}}, [%[[reg]]]
+  %addr1 = inttoptr i64 1 to i64 addrspace(1)*
+  store i64 42, i64* %addr0
+  %res = load i64, i64 addrspace(1)* %addr1
+  ret i64 %res
+}