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);
--- /dev/null
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py
+; RUN: llc -mtriple=nvptx64 < %s | FileCheck %s
+
+define i64 @test() nounwind readnone {
+; CHECK-LABEL: test(
+; CHECK: {
+; CHECK-NEXT: .reg .b64 %rd<4>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: mov.u64 %rd1, 1;
+; CHECK-NEXT: mov.u64 %rd2, 42;
+; CHECK-NEXT: st.u64 [%rd1], %rd2;
+; CHECK-NEXT: ld.global.u64 %rd3, [%rd1];
+; CHECK-NEXT: st.param.b64 [func_retval0+0], %rd3;
+; CHECK-NEXT: ret;
+ %addr0 = inttoptr i64 1 to ptr
+ %addr1 = inttoptr i64 1 to ptr addrspace(1)
+ store i64 42, ptr %addr0
+ %res = load i64, ptr addrspace(1) %addr1
+ ret i64 %res
+}