UNSPECV_CAS
UNSPECV_CAS_LOCAL
UNSPECV_XCHG
+ UNSPECV_ST
UNSPECV_BARSYNC
UNSPECV_WARPSYNC
UNSPECV_UNIFORM_WARP_CHECK
}
if (TARGET_SM70)
- /* Fall back to expand_atomic_store. */
- FAIL;
+ {
+ emit_insn (gen_nvptx_atomic_store<mode> (operands[0], operands[1],
+ operands[2]));
+ DONE;
+ }
bool maybe_shared_p = nvptx_mem_maybe_shared_p (operands[0]);
if (!maybe_shared_p)
DONE;
})
+(define_insn "nvptx_atomic_store<mode>"
+ [(set (match_operand:SDIM 0 "memory_operand" "+m") ;; memory
+ (unspec_volatile:SDIM
+ [(match_operand:SDIM 1 "nvptx_nonmemory_operand" "Ri") ;; input
+ (match_operand:SI 2 "const_int_operand")] ;; model
+ UNSPECV_ST))]
+ "TARGET_SM70"
+ {
+ const char *t
+ = "%.\tst%A0.b%T0\t%0, %1;";
+ return nvptx_output_atomic_insn (t, operands, 0, 2);
+ }
+ [(set_attr "atomic" "true")])
+
(define_insn "atomic_fetch_add<mode>"
[(set (match_operand:SDIM 1 "memory_operand" "+m")
(unspec_volatile:SDIM
--- /dev/null
+/* Test the atomic store expansion for sm > sm_6x targets,
+ shared state space. */
+
+/* { dg-do compile } */
+/* { dg-options "-misa=sm_75" } */
+
+enum memmodel
+{
+ MEMMODEL_SEQ_CST = 5
+};
+
+unsigned int u32 __attribute__((shared));
+unsigned long long int u64 __attribute__((shared));
+
+int
+main()
+{
+ __atomic_store_n (&u32, 0, MEMMODEL_SEQ_CST);
+ __atomic_store_n (&u64, 0, MEMMODEL_SEQ_CST);
+
+ return 0;
+}
+
+/* { dg-final { scan-assembler-times "st.shared.b32" 1 } } */
+/* { dg-final { scan-assembler-times "st.shared.b64" 1 } } */
+/* { dg-final { scan-assembler-times "membar.cta" 4 } } */