From 4aa8473543e5f3e7e8942d2a638eea3bc8c6c89e Mon Sep 17 00:00:00 2001
From: Henry Jin <hjin@nas.nasa.gov>
Date: Wed, 4 Feb 2015 23:01:13 -0800
Subject: [PATCH] update description and code for examples async_target.2[cf]

---
 Examples_async_target.tex         | 19 ++++++++++++++++++-
 Title_Page.tex                    |  4 +++-
 sources/Example_async_target.2c.c |  7 ++++---
 sources/Example_async_target.2f.f |  4 ++--
 4 files changed, 27 insertions(+), 7 deletions(-)

diff --git a/Examples_async_target.tex b/Examples_async_target.tex
index 04c79cd..9f7915c 100644
--- a/Examples_async_target.tex
+++ b/Examples_async_target.tex
@@ -24,7 +24,24 @@ accessed.
 
 \cexample{async_target}{2c}
 
-The Fortran example uses allocatable arrays for dynamic memory on the device. 
+The Fortran example below is similar to the C version above. Instead of pointers, though, it uses
+the convenience of Fortran allocatable arrays on the device. An allocatable array has the
+same behavior in a \code{map} clause as a C pointer, in this case.
+
+If there is no shape specified for an allocatable array in a \code{map} clause, only the array descriptor
+(also called a dope vector) is mapped. That is, device space is created for the descriptor, and it
+is initially populated with host values. In this case, the \plc{v1} and \plc{v2} arrays will be in a
+non-associated state on the device. When space for \plc{v1} and \plc{v2} is allocated on the device
+the addresses to the space will be included in their descriptors.
+
+At the end of the device region, the descriptor (of an unshaped specification of an allocatable
+array in a \code{map} clause) is returned with the raw device address of the allocated space.
+The content of the array is not returned. In the example the data in arrays \plc{v1} and \plc{v2}
+are not returned. In the second target directive, the \plc{v1} and \plc{v2} descriptors are
+re-created on the device with the descriptive information; and references to the
+vectors point to the correct local storage, of the space that was not freed in the first target
+directive.  At the end of the \code{target} region the data in array \plc{p} is copied back
+to the host, since \plc{p} is not an allocatable array.
 
 \fexample{async_target}{2f}
 
diff --git a/Title_Page.tex b/Title_Page.tex
index a6d1f58..bcd08eb 100644
--- a/Title_Page.tex
+++ b/Title_Page.tex
@@ -17,7 +17,7 @@
 
       \vspace{1.0in}
 
-      \textbf{Version 4.0.2 -- February, 2015}
+      \textbf{Version 4.0.2rv1 -- February, 2015}
     \end{center}
     \end{adjustwidth}
 
@@ -38,5 +38,7 @@ permission of OpenMP Architecture Review Board.\end{adjustwidth}
 \thispagestyle{empty}
 \phantom{a}
 \emph{This page intentionally left blank}
+
+This working version enacted the following tickets: 299, 342, and a few other editorial changes.
 \vfill
 
diff --git a/sources/Example_async_target.2c.c b/sources/Example_async_target.2c.c
index 5b65b7a..0662aa1 100644
--- a/sources/Example_async_target.2c.c
+++ b/sources/Example_async_target.2c.c
@@ -12,10 +12,11 @@ extern void init(float *, float *, int);
 #pragma omp end declare target
 extern void foo();
 extern void output(float *, int);
-void vec_mult(float *p, float *v1, float *v2, int N, int dev)
+void vec_mult(float *p, int N, int dev)
 {
+   float *v1, *v2;
    int i;
-   #pragma omp task depend(out: v1, v2)
+   #pragma omp task shared(v1, v2) depend(out: v1, v2)
    #pragma omp target device(dev) map(v1, v2)
    {
        // check whether on device dev
@@ -26,7 +27,7 @@ void vec_mult(float *p, float *v1, float *v2, int N, int dev)
        init(v1, v2, N);
    }
    foo(); // execute other work asychronously
-   #pragma omp task depend(in: v1, v2)
+   #pragma omp task shared(v1, v2, p) depend(in: v1, v2)
    #pragma omp target device(dev) map(to: v1, v2) map(from: p[0:N])
    {
        // check whether on device dev
diff --git a/sources/Example_async_target.2f.f b/sources/Example_async_target.2f.f
index 3953ece..0543d65 100644
--- a/sources/Example_async_target.2f.f
+++ b/sources/Example_async_target.2f.f
@@ -10,7 +10,7 @@
    integer ::  i, idev
    !$omp declare target (init)
 
-   !$omp task depend(out: v1,v2)
+   !$omp task shared(v1,v2) depend(out: v1,v2)
       !$omp target device(idev) map(v1,v2)
          if( omp_is_initial_device() ) &
             stop "not executing on target device"
@@ -21,7 +21,7 @@
 
    call foo()  ! execute other work asychronously
 
-   !$omp task depend(in: v1,v2)
+   !$omp task shared(v1,v2,p) depend(in: v1,v2)
       !$omp target device(idev) map(to: v1,v2) map(from: p)
          if( omp_is_initial_device() ) &
             stop "not executing on target device"