[Openmp-commits] [PATCH] D62397: [OPENMP][NVPTX]Relax flush directive.
Alexey Bataev via Phabricator via Openmp-commits
openmp-commits at lists.llvm.org
Wed Jun 19 08:17:27 PDT 2019
ABataev updated this revision to Diff 205599.
ABataev added a comment.
Address comments.
Repository:
rOMP OpenMP
CHANGES SINCE LAST ACTION
https://reviews.llvm.org/D62397/new/
https://reviews.llvm.org/D62397
Files:
libomptarget/deviceRTLs/nvptx/src/sync.cu
libomptarget/deviceRTLs/nvptx/test/parallel/flush.c
Index: libomptarget/deviceRTLs/nvptx/test/parallel/flush.c
===================================================================
--- /dev/null
+++ libomptarget/deviceRTLs/nvptx/test/parallel/flush.c
@@ -0,0 +1,33 @@
+// RUN: %compile-run-and-check
+
+#include <omp.h>
+#include <stdio.h>
+
+int main(int argc, char *argv[]) {
+ int data, flag = 0;
+#pragma omp target parallel num_threads(64) map(tofrom : data, flag)
+ {
+ if (omp_get_thread_num() == 0) {
+ /* Write to the data buffer that will be read by thread */
+ data = 42;
+/* Flush data to thread 1 and strictly order the write to data
+ relative to the write to the flag */
+#pragma omp flush(flag, data)
+ /* Set flag to release thread 1 */
+ flag = 1;
+/* Flush flag to ensure that thread 1 sees the change */
+#pragma omp flush(flag)
+ } else if (omp_get_thread_num() == 32) {
+/* Loop until we see the update to the flag */
+#pragma omp flush(flag, data)
+ while (flag < 1) {
+#pragma omp flush(flag, data)
+ }
+#pragma omp flush(flag, data)
+ }
+ }
+ // CHECK: data=42.
+ /* Value of data will be 42, value of flag still undefined */
+ printf("data=%d.\n", data);
+ return 0;
+}
Index: libomptarget/deviceRTLs/nvptx/src/sync.cu
===================================================================
--- libomptarget/deviceRTLs/nvptx/src/sync.cu
+++ libomptarget/deviceRTLs/nvptx/src/sync.cu
@@ -130,7 +130,7 @@
EXTERN void __kmpc_flush(kmp_Ident *loc) {
PRINT0(LD_IO, "call kmpc_flush\n");
- __threadfence_system();
+ __threadfence();
}
////////////////////////////////////////////////////////////////////////////////
-------------- next part --------------
A non-text attachment was scrubbed...
Name: D62397.205599.patch
Type: text/x-patch
Size: 1638 bytes
Desc: not available
URL: <http://lists.llvm.org/pipermail/openmp-commits/attachments/20190619/5894c804/attachment.bin>
More information about the Openmp-commits
mailing list