From 9fd75406f3f9643d69f8785063283d85a286b44a Mon Sep 17 00:00:00 2001
From: Kent Milfeld <milfeld@tacc.utexas.edu>
Date: Sat, 21 Oct 2023 15:10:55 -0500
Subject: [PATCH] Updated timer.hpp and timer.f90 to be consistent. Fleshed out
 time_add gpu exercise.

---
 OpenMP_cpu/adv_intro/2_timers/README_timer.md | 35 +++++----
 .../adv_intro/2_timers/solutions/ans_output   | 14 +++-
 .../2_timers/solutions/ans_time_print.cpp     |  4 +-
 .../2_timers/solutions/ans_time_print.f90     |  6 +-
 OpenMP_cpu/adv_intro/2_timers/time_print.cpp  |  2 +-
 OpenMP_cpu/adv_intro/2_timers/time_print.f90  |  2 +-
 OpenMP_cpu/adv_intro/2_timers/timer.f90       | 18 ++---
 OpenMP_cpu/adv_intro/2_timers/timer.hpp       | 29 +++----
 OpenMP_gpu/adv/time_add/README_time_add.md    | 47 +++++++++--
 OpenMP_gpu/adv/time_add/add.F90               | 66 ++++++++++++++++
 OpenMP_gpu/adv/time_add/add.cpp               | 72 ++++++++---------
 OpenMP_gpu/adv/time_add/timer.f90             | 78 +++++++++++++++++++
 OpenMP_gpu/adv/time_add/timer.hpp             | 29 +++----
 OpenMP_gpu/basic/worksharing/README.md        | 58 +++++++++-----
 .../solutions/Intel_needs_mcmodel_error       | 16 ++++
 .../solutions/ans_opt_worksharing.cpp         | 36 +++++++++
 .../solutions/ans_opt_worksharing.f90         | 40 ++++++++++
 .../worksharing/solutions/ans_worksharing.cpp |  8 +-
 .../worksharing/solutions/ans_worksharing.f90 | 34 ++++++++
 OpenMP_gpu/basic/worksharing/timer.f90        | 78 +++++++++++++++++++
 OpenMP_gpu/basic/worksharing/timer.hpp        | 30 +++----
 OpenMP_gpu/basic/worksharing/worksharing.f90  | 33 ++++++++
 22 files changed, 598 insertions(+), 137 deletions(-)
 create mode 100644 OpenMP_gpu/adv/time_add/add.F90
 create mode 100644 OpenMP_gpu/adv/time_add/timer.f90
 create mode 100644 OpenMP_gpu/basic/worksharing/solutions/Intel_needs_mcmodel_error
 create mode 100644 OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.cpp
 create mode 100644 OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.f90
 create mode 100644 OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.f90
 create mode 100644 OpenMP_gpu/basic/worksharing/timer.f90
 create mode 100644 OpenMP_gpu/basic/worksharing/worksharing.f90

diff --git a/OpenMP_cpu/adv_intro/2_timers/README_timer.md b/OpenMP_cpu/adv_intro/2_timers/README_timer.md
index ecc5c05..60c4a52 100644
--- a/OpenMP_cpu/adv_intro/2_timers/README_timer.md
+++ b/OpenMP_cpu/adv_intro/2_timers/README_timer.md
@@ -1,47 +1,56 @@
-# Class Timer  -- collects time measurements and reports them later.
+# Class Timer_Collector  -- collects time measurements and reports them later.
 
 ## BACKGROUND
 Fortran and CPP Timer classes have been created for measuring performance
 of code blocks demarked by start() and stop() class member functions.
 For each block a literal character string argument for start("my label")
 is used as a label when the times are reported with print().
-
-See timer.hpp (CPP) and timer.f90 (F90)
-
+```
+      Declare Timer_collector as timer
+               CPP                                F90
+      timer.start("my label")             call timer%start("my label")
+      <something to be timed>             <something to be timed> 
+      timer.stop()                        call timer%stop()
+      ...                                 ...
+      timer.print()                       call timer%print()
+
+      See timer.hpp (CPP) and timer.f90 (F90)
+
+```
 ## Exercises
 
 1.) 
-    Look over time_print(.cpp|.f90).  The todo's guide you
+    Look over time_print(.cpp|.f90).  The TODOs guide you
     through instrumenting the code for timing 
 
     a sleep routine, and
 
-    -- That's just to make sure the timer is working.---
+    -- That is just to make sure the timer is working.---
 
     a write to stdout (CPP: printf, F90: print*,)
     
     TODO 1:
-    Include the Timer class file (CPP: timer.hpp, F90: timer.f90)
+    Include the Timer_collector class file (CPP: timer.hpp, F90: timer.f90)
     at the beginning of the code.
 
     TODO 2:
     Instantiate the class in main as timer:
-    CPP:  Just use "Timer time;"
+    CPP:  Just use "Timer_Collector timer;"
     F90:  Use the module defined in timer.f90 (mod_timer)
-          and define timer as type(cls_timer).
+          and define "timer" as type(Timer_Collector).
 
     TODO 3:
     Time the sleep and print statements:
-    Put      time.start("<label>") and      time.stop() before and after
-    Put call time%start("<label>") and call time%stop() before and after
+    CPP: Put      time.start("my label") and      time.stop() before and after
+    F90: Put call time%start("my label") and call time%stop() before and after
 
     TODO 4:
     Report the times at the end.
 
     Is the timer reporting reasonable numbers?
     What are the units?
+    How accurate is the timer?
     How long did the print statement take?
-    What is the resolution of the timer.
+    (Note, the print time is just to write to an io buffer.)
     How much time is use in calling the timer?
     How would you measure that?
-    (TODO, fortran needs to report more digits)
diff --git a/OpenMP_cpu/adv_intro/2_timers/solutions/ans_output b/OpenMP_cpu/adv_intro/2_timers/solutions/ans_output
index a8f2029..91ea36b 100644
--- a/OpenMP_cpu/adv_intro/2_timers/solutions/ans_output
+++ b/OpenMP_cpu/adv_intro/2_timers/solutions/ans_output
@@ -1,11 +1,19 @@
 CPP code
 
+Ft1 solutions $ ./a.out
  Hello Helsinki.
 
 Action           ::    time/s     Time resolution = 1e-06
 ------
- 1 sec sleep  ::    1.00009
- printf io    ::    5.7e-05
-
+ 1 sec sleep  ::     1.000986
+ printf io    ::     0.000023
 
 Fortran code
+
+  Hello Helsinki.
+ 
+Action           ::   time/s     Time resolution =  1.0E-06
+------
+ 1 sec sleep     ::     1.000825
+ printf io       ::     0.000171
+
diff --git a/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.cpp b/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.cpp
index f0ae389..7e526fc 100644
--- a/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.cpp
+++ b/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.cpp
@@ -5,8 +5,8 @@
 #include "timer.hpp"
 int main(){
 
-// TODO 1a: Instantiate a Timer class as timer.
-   Timer timer;
+// TODO 1a: Instantiate  a Timer_Collector class as timer.
+   Timer_Collector timer;
 
 // TODO 1b: Time a 1 second sleep
    timer.start(" 1 sec sleep ");
diff --git a/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.f90 b/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.f90
index c4d48e9..f3ed177 100644
--- a/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.f90
+++ b/OpenMP_cpu/adv_intro/2_timers/solutions/ans_time_print.f90
@@ -4,9 +4,9 @@
 program main
 
 !! TODO 1a: Use the module in timer.f90 
-!!          and instantiate a type(cls_timer) class
-   use mod_timer
-   type(cls_timer) timer
+!!          and instantiate Timer_Collector (type) class as timer
+   use mod_Timer
+   type(Timer_Collector) timer
 
 !! TODO 1b: Time a 1 second sleep
    call timer%start(" 1 sec sleep ");
diff --git a/OpenMP_cpu/adv_intro/2_timers/time_print.cpp b/OpenMP_cpu/adv_intro/2_timers/time_print.cpp
index 38dd311..1d2b075 100644
--- a/OpenMP_cpu/adv_intro/2_timers/time_print.cpp
+++ b/OpenMP_cpu/adv_intro/2_timers/time_print.cpp
@@ -5,7 +5,7 @@
 // #include ...
 int main(){
 
-// TODO 1a: Instantiate a Timer class as timer.
+// TODO 1a: Instantiate a Timer_Collector class as timer.
 // 
 
 // TODO 1b: Time a 1 second sleep
diff --git a/OpenMP_cpu/adv_intro/2_timers/time_print.f90 b/OpenMP_cpu/adv_intro/2_timers/time_print.f90
index a10e055..e53502d 100644
--- a/OpenMP_cpu/adv_intro/2_timers/time_print.f90
+++ b/OpenMP_cpu/adv_intro/2_timers/time_print.f90
@@ -4,7 +4,7 @@
 program main
 
 !! TODO 1a: Use the module in timer.f90 
-!!          and instantiate a type(cls_timer) class
+!!          and instantiate Timer_Collector (type) class as timer
   !use  ...
   !type(...) timer
 
diff --git a/OpenMP_cpu/adv_intro/2_timers/timer.f90 b/OpenMP_cpu/adv_intro/2_timers/timer.f90
index cae8a3b..70ad83d 100644
--- a/OpenMP_cpu/adv_intro/2_timers/timer.f90
+++ b/OpenMP_cpu/adv_intro/2_timers/timer.f90
@@ -4,7 +4,7 @@ module mod_timer
    private
    integer, parameter     :: m = 20
    
-   type, public :: cls_timer
+   type, public :: Timer_Collector
    
    private
    integer                         :: n = 0
@@ -19,13 +19,13 @@ module mod_timer
    procedure, public :: stop
    procedure, public :: print
    
-   endtype cls_timer
+   endtype Timer_Collector
    
    contains
    
    subroutine reset(this)
    
-   class(cls_timer) :: this
+   class(Timer_Collector) :: this
    this%n  = 0
    this%it = -1
    this%c  = 'undef'
@@ -35,7 +35,7 @@ subroutine reset(this)
    
    subroutine start(this, c)
    
-   class(cls_timer) :: this
+   class(Timer_Collector) :: this
    character(len=*) :: c
    character(len=16) :: c16
    
@@ -54,7 +54,7 @@ subroutine start(this, c)
    
    subroutine stop(this)
    
-   class(cls_timer) :: this
+   class(Timer_Collector) :: this
    
    call system_clock(this%it(2,this%n))
    
@@ -63,14 +63,14 @@ subroutine stop(this)
    
    subroutine print(this)
    
-   class(cls_timer) :: this
+   class(Timer_Collector) :: this
    
    write (*,*)
-   write (*,'(a,es7.1)') 'Action           ::   time/s     Time resolution = ', 1./real(this%itr)
+   write (*,'(a,es8.1)') 'Action           ::   time/s     Time resolution = ', 1./real(this%itr)
    write (*,'(a)')       '------'
    do i=1, this%n
-     write (*,'(a,a, f7.3)') this%c(i), ' :: ',               &
-                               (real(this%it(2,i) - this%it(1,i))) / real(this%itr)
+     write (*,'(a,a, f12.6)') this%c(i), ' :: ',               &  !TODO change .6 to .precision
+                              (real(this%it(2,i) - this%it(1,i))) / real(this%itr)
    enddo
    
    end subroutine
diff --git a/OpenMP_cpu/adv_intro/2_timers/timer.hpp b/OpenMP_cpu/adv_intro/2_timers/timer.hpp
index 2653b38..659b49e 100644
--- a/OpenMP_cpu/adv_intro/2_timers/timer.hpp
+++ b/OpenMP_cpu/adv_intro/2_timers/timer.hpp
@@ -2,44 +2,47 @@
 #define SIMPLE_TIMER_H
 
 #include <iostream>
+#include <iomanip>    //setprecision
 #include <cmath>
 #include <string>
 #include <sys/time.h>
+using namespace std;
 
 /* simple timer that stores multiple times 
    and can print them afterwards.
    Does not support nested timing calls. 
 */
-class Timer
+class Timer_Collector
 {
 public:
-    Timer(): n(0) { }
-    void start(std::string label) 
+    Timer_Collector(): n(0) { }
+    void start(string label) 
     { 
         if (n < 20) { labels[n] = label; gettimeofday(&times[2*n], NULL); }
-        else { std::cerr << "No more timers, " << label << " will not be timed." << std::endl; }
+        else { cerr << "No more timers, " << label << " will not be timed." << endl; }
     }
 
     void stop() { gettimeofday(&times[2*n+1], NULL); n++;}
     void reset() { n=0; }
     void print();
 private:
-    std::string labels[20];
+    string labels[20];
     timeval      times[40];
     int n;
 };
 
-void Timer::print()
+void Timer_Collector::print()
 {
-    std::cout << std::endl;
-    std::cout << "Action           ::    time/s     Time resolution = " << 1.f/(float)CLOCKS_PER_SEC << std::endl;
-    std::cout << "------" << std::endl;
+    cout << endl;
+    cout << "Action           ::    time/s     Time resolution = " << 1.f/(float)CLOCKS_PER_SEC << endl;
+    cout << "------" << endl;
     for (int i=0; i < n; ++i)
     {
-        time_t seconds = times[2*i+1].tv_sec - times[2*i+0].tv_sec;
-        suseconds_t ms = times[2*i+1].tv_usec - times[2*i+0].tv_usec;
-        if (ms < 0) { ms += 1000000; seconds--; }
-        std::cout << labels[i] << " ::    " << (float)seconds + ms/1000000.f << std::endl;
+        time_t seconds = times[2*i+1].tv_sec  - times[2*i+0].tv_sec;
+        suseconds_t us = times[2*i+1].tv_usec - times[2*i+0].tv_usec;
+        if (us < 0) { us += 1000000; seconds--; }
+        double secs = (double)seconds + us/1000000.0f;
+        cout << labels[i] << " :: " << setiosflags(ios::fixed)<<setprecision(6)<< setw(12)<<secs<<endl;
     }
 }
 #endif
diff --git a/OpenMP_gpu/adv/time_add/README_time_add.md b/OpenMP_gpu/adv/time_add/README_time_add.md
index aa093e9..b57cf2b 100644
--- a/OpenMP_gpu/adv/time_add/README_time_add.md
+++ b/OpenMP_gpu/adv/time_add/README_time_add.md
@@ -1,10 +1,45 @@
-# This exercise is incomplete.
+# Time Offload kernel for vector add, and data motion
+
+## Background
+
+By putting timers around the "target data" directive"
+and immediately after the target data region begind,,
+the "target" construct, and the termination of the target
+data region, one can measure the time required to 
+
+```
+   (allocate device data and) move data "to" the device, 
+
+   execute and kernel, 
+
+   and move data "from" the device (and deallocate).
+```
 
 ## Exercise
-Only C/C++ code is available at this time
-1.) Experiment with the target ... directives 
-    executing the addition.
-    Some different setting are in comments.
+    STATUS:  Makefile works for CPP only at this time.
 
+1.) Add the start and stop timers around the appropriate
+    areas, and run with the default (present) "target"
+    construct to offload the add function.  Make sue it works.
+    See previous examples (basic/worksharing) for including a timer:
+    TODO 1:
     Note, the data motions due to the target data
-    construct are measured separately.
+    construct are measured separately at the beginning and end
+    of the target data region.
+    
+    TODO 2:  2a target ENCLOSING function, 2b constructs IN function
+    Change 1<<25 to 1<<28 for C/C++; 2**25 to 2**28 for Fortran for timing.
+
+    Note: For the "add" procedure, the target and the target teams can be
+    hoisted outside the function.  The different forms are labeled with
+    F1, F2, F3, and F4.  Experiment with (time) the differents forms,
+    and determine which one works best with defaults, and then experiment
+    with the number of teams and thread limit for the optimal version.
+    Which one has the highest performance.
+
+    TODO 3:
+    Now try running the Loop version (which doesn't
+    call the function) L1, L2, L3, with various different 
+    "target" directives (number of teams and thread limit).
+
+    Which one has the highest performance.
diff --git a/OpenMP_gpu/adv/time_add/add.F90 b/OpenMP_gpu/adv/time_add/add.F90
new file mode 100644
index 0000000..465967a
--- /dev/null
+++ b/OpenMP_gpu/adv/time_add/add.F90
@@ -0,0 +1,66 @@
+include "timer.f90"
+
+subroutine add(n, x, y)
+  implicit none
+  integer :: n,i
+  real    :: x(n),y(n)
+
+    !! TODO 2b:
+!F1 !$omp target teams distribute parallel for num_teams(108) num_threads(256) !!1 .4s
+!F2 !$omp        teams distribute parallel for num_teams(108) num_threads(256) !!2 mins
+!F3 !$omp              distribute parallel for                num_threads(256) !!2 .1s
+    !$omp              distribute parallel for !!             num_threads(256) !!2 .1s
+    do i=1,n;  y(i)=x(i)+y(i);  enddo
+
+end subroutine 
+
+program main
+  use mod_Timer
+  implicit none
+
+  type(Timer_Collector) timer
+
+ !integer,parameter :: N=2**28 !! use 28 for timing, 1024*1024*2**8
+  integer,parameter :: N=2**25 !! use 25 for code validation,  1024*1024*2**5
+  real    :: x(N), y(N), error
+  integer :: i
+
+
+  do i=1,n;  x(i)=1.0e0; y(i)=1.0e0; enddo  !init x and y on host
+
+  call timer%start(" Data TO   & Alloc  "); !! TODO 1
+    !$omp target data map(tofrom:y) map(to:x) !allocate and move data
+  call timer%stop();                        !! TODO 1
+ 
+  call timer%start(" Add on GPU         "); !! TODO 1
+
+    !! TODO 3:
+    !!    Time Loop on Host
+    !!#pragma omp target teams distribute parallel for num_teams(108) num_threads(256)
+    !!do i=1,n;  y(i)=x(i)+y(i);  enddo
+
+    !! TODO 2a:
+    !!    Time function call with "target" "target teams and num_teams/thread_limit clauses
+    !F1 -- no construct here                      !!1 .4s
+    !F2 #pragma !$omp target                      !!3  mins
+    !F3 #pragma !$omp target teams num_teams(108) !!2 .1s
+    !F4 #pragma !$omp target teams thread_limit(256) !! num_teams(108) !!2 .008s
+    !$omp target teams thread_limit(512) !! num_teams(108) !!2 
+    call add(N,x,y);
+
+  call timer%stop();                        !! TODO 1
+
+  call timer%start(" Data FROM & Dealloc"); !! TODO 1
+    !$omp end target data 
+  call timer%stop();                        !! TODO 1
+
+  do i=1,N
+     error = max( 0.0e0, abs(y(i)-2.0e0))
+     if ( error > 0.0000001 ) &
+        print*, " error e10-6 exceeded at ",i," val= ",error
+  end do
+  print*, "Max error = ", error, " for count of ", N
+
+  call timer%print();                       !! TODO 1
+
+end program main
diff --git a/OpenMP_gpu/adv/time_add/add.cpp b/OpenMP_gpu/adv/time_add/add.cpp
index b759b1d..33dd1b6 100644
--- a/OpenMP_gpu/adv/time_add/add.cpp
+++ b/OpenMP_gpu/adv/time_add/add.cpp
@@ -5,58 +5,60 @@ using namespace std;
 
 void add(int n, float *x, float *y)
  { 
-//1 #pragma omp target teams distribute parallel for num_teams(108) num_threads(256) //1 .4s
-//2 #pragma omp        teams distribute parallel for num_teams(108) num_threads(256) //2 mins
-//3 #pragma omp              distribute parallel for                num_threads(256) //2 .1s
-    #pragma omp              distribute parallel for //             num_threads(256) //2 .1s
+// TODO 2b
+//F1 #pragma omp target teams distribute parallel for num_teams(108) num_threads(256) //1 .4s
+//F2 #pragma omp        teams distribute parallel for num_teams(108) num_threads(256) //2 mins
+//F3 #pragma omp              distribute parallel for                num_threads(256) //2 .1s
+//F4 #pragma omp              distribute parallel for //             num_threads(256) //2 .1s
    for (int i = 0; i < n; i++) y[i] = x[i] + y[i]; 
  }
 
 int main(void)
 {
-  Timer timer;
-  int N = 1<<28; // 1M elements
-//int N = 1<<25; // 1M elements
-  float x[N], y[N]; 
-
-  timer.start("HOST init    ");
-  // initialize x and y arrays on the host
-  for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; }
- #pragma omp target data map(tofrom:y) map(to:x)
-{
-  timer.stop();
- 
+  Timer_Collector timer;
+//int N = 1<<28; //  use 2**28 for timings 
+  int N = 1<<25; //  use 2**25 for validating code
+  float x[N], y[N], y_value, maxError = 0.0f;
+
+  for (int i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } //init x & Y on host
+  y_value=x[0]+y[0];
+
+timer.start(" Alloc and Data TO  ");        // TODO 1 :
 
-//   for (int i = 0; i < N; i++) y[i] = x[i] + y[i];
+  #pragma omp target data map(tofrom:y) map(to:x)
+  {
+timer.stop();                               // TODO 1 :
 
+timer.start(" Add on GPU         ");        // TODO 1 :
 
-  timer.start(" GPU  add    ");
-//#pragma omp target teams distribute parallel for num_teams(108) num_threads(256)
-//for (int i = 0; i < N; i++) y[i] = x[i] + y[i];
+  // TODO 
+  //#pragma omp target teams distribute parallel for num_teams(108) num_threads(256)
+  //for (int i = 0; i < N; i++) y[i] = x[i] + y[i];
 
-//1                                         //1 .4s
-//2 #pragma omp target                      //3  mins
-//3 #pragma omp target teams num_teams(108) //2 .1s
-//4 #pragma omp target teams thread_limit(256) // num_teams(108) //2 .008s
-    #pragma omp target teams thread_limit(512) // num_teams(108) //2 
+  // TODO 2b
+  //F1                                         //1 .4s
+  //F2 #pragma omp target                      //3  mins
+  //F3 #pragma omp target teams num_teams(108) //2 .1s
+  //F4 #pragma omp target teams thread_limit(256) // num_teams(108) //2 .008s
   add(N,x,y);
  
-  timer.stop();
+timer.stop();                               // TODO 1 :
 
 
 
-  timer.start("HOST check   ");
-}
+timer.start(" Data FROM & Dealloc");        // TODO 1 :
+  }
+timer.stop();                               // TODO 1 :
+
   // Check for errors (all values should be 3.0f)
-  float maxError = 0.0f;
   for (int i = 0; i < N; i++){
-    maxError = fmax(maxError, fabs(y[i]-3.0f));
-    if(maxError > 0.0000001) {cout << i << maxError << " at i=" <<  i <<  endl; }
+    maxError = fmax(maxError, fabs(y[i]-2.0f));
+    if(maxError > 0.0000001) {cout << maxError<< " error exceeds e10-6 at i=" << i << endl;}
   }
-  std::cout << "Max error: " << maxError << std::endl;
-  std::cout << " for count = " << N << std::endl;
-  timer.stop();
-  timer.print();
+  std::cout << " Max error:  " << maxError << std::endl;
+  std::cout << " for count = " << N        << std::endl;
+
+  timer.print();                            // TODO 1 :
 
   return 0;
 }
diff --git a/OpenMP_gpu/adv/time_add/timer.f90 b/OpenMP_gpu/adv/time_add/timer.f90
new file mode 100644
index 0000000..70ad83d
--- /dev/null
+++ b/OpenMP_gpu/adv/time_add/timer.f90
@@ -0,0 +1,78 @@
+  module mod_timer
+   ! Developed by Lars Koesterke @ TACC 
+   
+   private
+   integer, parameter     :: m = 20
+   
+   type, public :: Timer_Collector
+   
+   private
+   integer                         :: n = 0
+   integer(8), dimension(2,m)      :: it
+   integer(8)                      :: itr
+   character(len=16), dimension(m) :: c
+   
+   contains
+   
+   procedure, public :: reset
+   procedure, public :: start
+   procedure, public :: stop
+   procedure, public :: print
+   
+   endtype Timer_Collector
+   
+   contains
+   
+   subroutine reset(this)
+   
+   class(Timer_Collector) :: this
+   this%n  = 0
+   this%it = -1
+   this%c  = 'undef'
+   
+   end subroutine
+   
+   
+   subroutine start(this, c)
+   
+   class(Timer_Collector) :: this
+   character(len=*) :: c
+   character(len=16) :: c16
+   
+   this%n = this%n + 1
+   if (this%n > m) then
+     write (*,*) 'No more timers available'
+     stop 'Fatal ERROR in method START of class CLS_TIMER'
+   endif
+   call system_clock(this%it(1,this%n), this%itr)
+   ! to get PGI 11 going
+   c16 = c
+   this%c(this%n) = c16
+   
+   end subroutine
+   
+   
+   subroutine stop(this)
+   
+   class(Timer_Collector) :: this
+   
+   call system_clock(this%it(2,this%n))
+   
+   end subroutine
+   
+   
+   subroutine print(this)
+   
+   class(Timer_Collector) :: this
+   
+   write (*,*)
+   write (*,'(a,es8.1)') 'Action           ::   time/s     Time resolution = ', 1./real(this%itr)
+   write (*,'(a)')       '------'
+   do i=1, this%n
+     write (*,'(a,a, f12.6)') this%c(i), ' :: ',               &  !TODO change .6 to .precision
+                              (real(this%it(2,i) - this%it(1,i))) / real(this%itr)
+   enddo
+   
+   end subroutine
+   
+ end module mod_timer
diff --git a/OpenMP_gpu/adv/time_add/timer.hpp b/OpenMP_gpu/adv/time_add/timer.hpp
index 2653b38..659b49e 100644
--- a/OpenMP_gpu/adv/time_add/timer.hpp
+++ b/OpenMP_gpu/adv/time_add/timer.hpp
@@ -2,44 +2,47 @@
 #define SIMPLE_TIMER_H
 
 #include <iostream>
+#include <iomanip>    //setprecision
 #include <cmath>
 #include <string>
 #include <sys/time.h>
+using namespace std;
 
 /* simple timer that stores multiple times 
    and can print them afterwards.
    Does not support nested timing calls. 
 */
-class Timer
+class Timer_Collector
 {
 public:
-    Timer(): n(0) { }
-    void start(std::string label) 
+    Timer_Collector(): n(0) { }
+    void start(string label) 
     { 
         if (n < 20) { labels[n] = label; gettimeofday(&times[2*n], NULL); }
-        else { std::cerr << "No more timers, " << label << " will not be timed." << std::endl; }
+        else { cerr << "No more timers, " << label << " will not be timed." << endl; }
     }
 
     void stop() { gettimeofday(&times[2*n+1], NULL); n++;}
     void reset() { n=0; }
     void print();
 private:
-    std::string labels[20];
+    string labels[20];
     timeval      times[40];
     int n;
 };
 
-void Timer::print()
+void Timer_Collector::print()
 {
-    std::cout << std::endl;
-    std::cout << "Action           ::    time/s     Time resolution = " << 1.f/(float)CLOCKS_PER_SEC << std::endl;
-    std::cout << "------" << std::endl;
+    cout << endl;
+    cout << "Action           ::    time/s     Time resolution = " << 1.f/(float)CLOCKS_PER_SEC << endl;
+    cout << "------" << endl;
     for (int i=0; i < n; ++i)
     {
-        time_t seconds = times[2*i+1].tv_sec - times[2*i+0].tv_sec;
-        suseconds_t ms = times[2*i+1].tv_usec - times[2*i+0].tv_usec;
-        if (ms < 0) { ms += 1000000; seconds--; }
-        std::cout << labels[i] << " ::    " << (float)seconds + ms/1000000.f << std::endl;
+        time_t seconds = times[2*i+1].tv_sec  - times[2*i+0].tv_sec;
+        suseconds_t us = times[2*i+1].tv_usec - times[2*i+0].tv_usec;
+        if (us < 0) { us += 1000000; seconds--; }
+        double secs = (double)seconds + us/1000000.0f;
+        cout << labels[i] << " :: " << setiosflags(ios::fixed)<<setprecision(6)<< setw(12)<<secs<<endl;
     }
 }
 #endif
diff --git a/OpenMP_gpu/basic/worksharing/README.md b/OpenMP_gpu/basic/worksharing/README.md
index de8febc..6f0800f 100644
--- a/OpenMP_gpu/basic/worksharing/README.md
+++ b/OpenMP_gpu/basic/worksharing/README.md
@@ -1,62 +1,80 @@
 # Parallel Work with OpenMP offloading
 
-
-
-The target teams distribute parallel for (ttdpf) construct 
-(what a combined directive) parallelizes a for/do loop across 
+The "target teams distribute parallel for" (TTDPF) construct 
+(what a long combined directive!) parallelizes a for/do loop across 
 teams for threaded execution using a default number of teams and threads.
    o First approach is to just let the compiler decide the
      best combination of thread per team and the number of teams.
 
 1.) Insert the TTDPF directive to offload the axpy operation.
-    Fortran folks, just use the cpp code.
 
     Compile and run to make sure it work.
     Did you need a header file?  Why not.
 
+```
+    F90: Intel compiler needs " -mcmodel=medium " option.
+         ifort -mcmodel=medium  ...
+```
+
 
 2.) Now instrument the file with the included timer.
     a.) insert the declaration:
-           Timer timer;
-    b.) around the TTDPF region, start, stop, and print
+        CPP:  Just use "Timer_Collector timer;"
+        F90:  Use the module defined in timer.f90 (mod_timer)
+              and define "timer" as type(Timer_Collector).
+    b.) Around the TTDPF region, start, stop, and print
         the timer results:
-
-           timer.start("    Offload     ");
+```
+        CPP:
+        timer.start("    Offload     ");
            ...
         timer.stop();
         timer.print(); 
 
+        F90:
+        call timer%start("    Offload     ");
+           ...
+        call timer%stop();
+        call timer%print(); 
+```
+
     c.) Compile and run.
         How much time did it take?
-        Remember, moved data to the device (in a non optimal way)
+        Remember, data is being moved to the device (in a non-optimal way)
 
     d.) Now increase the size of the problem by 4x:
       
-        change 1<<29  to 1<<31
+        change 1<<28  to 1<<30
 
     e.) Compile and run
         Was the total cost 4x?
-        Most of the was from data motion.
-        We will look at that later.
+        Was most of the time used for data motion?
+        (How can you support your answer?
+        We will look at that later.)
+
+Include the following code if you want to check the results:
 
 ```
-Include the following code if youo want to check the results:
+CPP:
 
 //Check for rounding errors
 
 //Include following headers and namespace 
-//statement at beginning
+//        statement at beginning
 #include <cassert>
 #include <cmath>
 #include <iostream>
 using namespace std;
 
   //Insert after axpy calculation
-  float maxerr = 0.1f;
   for (int i=0; i<N; i++)
-  {
-    maxerr = fmax(maxerr, fabs(y[i]-3.0f));
-    assert(maxerr < 0.0000001); 
-  }
+    assert( (fabs(y[i]-3.0f)) < 0.000001);
+  cout << " Passed Validation "<< endl;
+
+F90:
+  do i = 1,N 
+    if( abs(y(i)-3.0e0) > 0.0000001 ) stop("Failed Validation");  
+  enddo
+  print*," Passed Validation"
 
 ```
diff --git a/OpenMP_gpu/basic/worksharing/solutions/Intel_needs_mcmodel_error b/OpenMP_gpu/basic/worksharing/solutions/Intel_needs_mcmodel_error
new file mode 100644
index 0000000..bb51d54
--- /dev/null
+++ b/OpenMP_gpu/basic/worksharing/solutions/Intel_needs_mcmodel_error
@@ -0,0 +1,16 @@
+iforolutions $ ifort ans_worksharing.f90 
+/tmp/ifortPGgpfh.o: in function `MAIN__':
+ans_worksharing.f90:(.text+0x5f): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+ans_worksharing.f90:(.text+0x6d): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+ans_worksharing.f90:(.text+0x96): relocation truncated to fit: R_X86_64_32S against `.bss'
+ans_worksharing.f90:(.text+0xb7): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+ans_worksharing.f90:(.text+0xc3): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+ans_worksharing.f90:(.text+0xce): relocation truncated to fit: R_X86_64_PC32 against `mod_timer_mp_start_$C16.0.3'
+ans_worksharing.f90:(.text+0xd5): relocation truncated to fit: R_X86_64_32S against `.bss'
+ans_worksharing.f90:(.text+0x117): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+ans_worksharing.f90:(.text+0x142): relocation truncated to fit: R_X86_64_32S against `.bss'
+ans_worksharing.f90:(.text+0x1a7): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+ans_worksharing.f90:(.text+0x2db): additional relocation overflows omitted from the output
+/opt/apps/gcc/8.3.0/bin/ld: failed to convert GOTPCREL relocation; relink with --no-relax ans_worksharing.f90
+
+
diff --git a/OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.cpp b/OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.cpp
new file mode 100644
index 0000000..efb94be
--- /dev/null
+++ b/OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.cpp
@@ -0,0 +1,36 @@
+#include <cassert>
+#include <cmath>
+#include <iostream>
+using namespace std;
+
+// TODO 2: include timer.hpp here
+#include "timer.hpp"
+int main(void)
+{
+//int     N = 1<<28;
+  int     N = 1<<30;
+  float x[N], y[N], a=2.0f; 
+
+  // TODO 2: Instantiate timer object here 
+  Timer_Collector timer;
+
+  // initialize x and y arrays on the host
+  for (int i = 0; i < N; i++) { x[ i] = 1.0f; y[ i] = 1.0f; }
+
+  // TODO 2: Start timer here
+  timer.start("    Offload     ");
+  // TODO 1: Offload with teams distribute parallel for
+  #pragma omp target teams distribute parallel for        //blk
+  { for (int i = 0; i < N; i++) y[i] = a*x[i] + y[i]; }
+
+  // TODO 2: Stop timer here and report resuls
+  timer.stop();
+  timer.print();
+
+  //Insert after axpy calculation
+  for (int i=0; i<N; i++)
+    assert( (fabs(y[i]-3.0f)) < 0.000001);
+  cout << " Passed Validation "<< endl;
+
+  return 0;
+}
diff --git a/OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.f90 b/OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.f90
new file mode 100644
index 0000000..8437ef9
--- /dev/null
+++ b/OpenMP_gpu/basic/worksharing/solutions/ans_opt_worksharing.f90
@@ -0,0 +1,40 @@
+! if Intel compile emits this error:
+!  *.f90:(.text+0x5f): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+! Compile with:
+! ifort -mcmodel=medium   *.f90
+
+  !! TODO 2: include  timer.f90 here
+include "timer.f90"
+
+program main
+  !! TODO 2: use mod_timer here
+  use mod_timer
+  implicit none
+ !integer, parameter :: N=1024*1024*2**8   !2**28
+  integer, parameter :: N=1024*1024*2**10  !2**30
+  real               :: x(N), y(N), a=2.0e0; 
+  integer            :: i
+
+  !! TODO 2: Instantiate timer object, Timer_collector, as timer.
+  type(Timer_Collector) timer
+
+  !! initialize x and y arrays on the host
+  do i = 1,N;  x(i) = 1.0e0; y(i) = 1.0e0;  enddo
+
+  !! TODO 2: Start timer here
+  call timer%start("    Offload     ");
+
+  !! TODO 1: Offload with teams distribute parallel for
+    !$omp target teams distribute parallel for        !!blk
+    do i = 1,N;  y(i) = a*x(i) + y(i); enddo
+
+  !! TODO 2: Stop timer here and report resuls
+  call timer%stop();
+  call timer%print();
+
+  do i = 1,N 
+     if( abs(y(i)-3.0e0) > 0.0000001) stop("Failed Validation")
+  enddo
+  print*," Passed Validation"
+
+end program main
diff --git a/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.cpp b/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.cpp
index cb32968..b4a5f18 100644
--- a/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.cpp
+++ b/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.cpp
@@ -3,11 +3,12 @@
 #include "timer.hpp"
 int main(void)
 {
-  int     N = 1<<31;
+  int     N = 1<<28;
+//int     N = 1<<30;
   float x[N], y[N], a=2.0f; 
 
-// TODO 2: Instantiate timer object here 
-  Timer timer;
+  // TODO 2: Instantiate timer object here 
+  Timer_Collector timer;
 
   // initialize x and y arrays on the host
   for (int i = 0; i < N; i++) { x[ i] = 1.0f; y[ i] = 1.0f; }
@@ -24,4 +25,3 @@ int main(void)
 
   return 0;
 }
-
diff --git a/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.f90 b/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.f90
new file mode 100644
index 0000000..16f961e
--- /dev/null
+++ b/OpenMP_gpu/basic/worksharing/solutions/ans_worksharing.f90
@@ -0,0 +1,34 @@
+  !! TODO 2: include  timer.f90 here
+include "timer.f90"
+
+program main
+  !! TODO 2: use mod_timer here
+  use mod_timer
+  implicit none
+  integer, parameter :: N=1024*1024*2**8   !2**28
+ !integer, parameter :: N=1024*1024*2**10  !2**30
+  real               :: x(N), y(N), a=2.0e0; 
+  integer            :: i
+
+  !! TODO 2: Instantiate timer object, Timer_collector, as timer.
+  type(Timer_Collector) timer
+
+  !! initialize x and y arrays on the host
+  do i = 1,N;  x(i) = 1.0e0; y(i) = 1.0e0;  enddo
+
+  !! TODO 2: Start timer here
+  call timer%start("    Offload     ");
+
+  !! TODO 1: Offload with teams distribute parallel for
+    !$omp target teams distribute parallel for        !!blk
+    do i = 1,N;  y(i) = a*x(i) + y(i); enddo
+
+  !! TODO 2: Stop timer here and report resuls
+  call timer%stop();
+  call timer%print();
+
+end program main
+! if Intel compile emits this error:
+!  *.f90:(.text+0x5f): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+! Compile with:
+! ifort -mcmodel=medium   *.f90
diff --git a/OpenMP_gpu/basic/worksharing/timer.f90 b/OpenMP_gpu/basic/worksharing/timer.f90
new file mode 100644
index 0000000..70ad83d
--- /dev/null
+++ b/OpenMP_gpu/basic/worksharing/timer.f90
@@ -0,0 +1,78 @@
+  module mod_timer
+   ! Developed by Lars Koesterke @ TACC 
+   
+   private
+   integer, parameter     :: m = 20
+   
+   type, public :: Timer_Collector
+   
+   private
+   integer                         :: n = 0
+   integer(8), dimension(2,m)      :: it
+   integer(8)                      :: itr
+   character(len=16), dimension(m) :: c
+   
+   contains
+   
+   procedure, public :: reset
+   procedure, public :: start
+   procedure, public :: stop
+   procedure, public :: print
+   
+   endtype Timer_Collector
+   
+   contains
+   
+   subroutine reset(this)
+   
+   class(Timer_Collector) :: this
+   this%n  = 0
+   this%it = -1
+   this%c  = 'undef'
+   
+   end subroutine
+   
+   
+   subroutine start(this, c)
+   
+   class(Timer_Collector) :: this
+   character(len=*) :: c
+   character(len=16) :: c16
+   
+   this%n = this%n + 1
+   if (this%n > m) then
+     write (*,*) 'No more timers available'
+     stop 'Fatal ERROR in method START of class CLS_TIMER'
+   endif
+   call system_clock(this%it(1,this%n), this%itr)
+   ! to get PGI 11 going
+   c16 = c
+   this%c(this%n) = c16
+   
+   end subroutine
+   
+   
+   subroutine stop(this)
+   
+   class(Timer_Collector) :: this
+   
+   call system_clock(this%it(2,this%n))
+   
+   end subroutine
+   
+   
+   subroutine print(this)
+   
+   class(Timer_Collector) :: this
+   
+   write (*,*)
+   write (*,'(a,es8.1)') 'Action           ::   time/s     Time resolution = ', 1./real(this%itr)
+   write (*,'(a)')       '------'
+   do i=1, this%n
+     write (*,'(a,a, f12.6)') this%c(i), ' :: ',               &  !TODO change .6 to .precision
+                              (real(this%it(2,i) - this%it(1,i))) / real(this%itr)
+   enddo
+   
+   end subroutine
+   
+ end module mod_timer
diff --git a/OpenMP_gpu/basic/worksharing/timer.hpp b/OpenMP_gpu/basic/worksharing/timer.hpp
index 3dd2270..659b49e 100644
--- a/OpenMP_gpu/basic/worksharing/timer.hpp
+++ b/OpenMP_gpu/basic/worksharing/timer.hpp
@@ -2,45 +2,47 @@
 #define SIMPLE_TIMER_H
 
 #include <iostream>
+#include <iomanip>    //setprecision
 #include <cmath>
 #include <string>
 #include <sys/time.h>
+using namespace std;
 
 /* simple timer that stores multiple times 
    and can print them afterwards.
    Does not support nested timing calls. 
 */
-class Timer
+class Timer_Collector
 {
 public:
-    Timer(): n(0) { }
-    void start(std::string label) 
+    Timer_Collector(): n(0) { }
+    void start(string label) 
     { 
         if (n < 20) { labels[n] = label; gettimeofday(&times[2*n], NULL); }
-        else { std::cerr << "No more timers, " << label << " will not be timed." << std::endl; }
+        else { cerr << "No more timers, " << label << " will not be timed." << endl; }
     }
 
     void stop() { gettimeofday(&times[2*n+1], NULL); n++;}
     void reset() { n=0; }
     void print();
 private:
-    std::string labels[20];
+    string labels[20];
     timeval      times[40];
     int n;
 };
 
-void Timer::print()
+void Timer_Collector::print()
 {
-    std::cout << std::endl;
-    std::cout << "    Action       ::    time/s     Time resolution = " << 1.f/(float)CLOCKS_PER_SEC << std::endl;
-    std::cout << "-----------------" << std::endl;
+    cout << endl;
+    cout << "Action           ::    time/s     Time resolution = " << 1.f/(float)CLOCKS_PER_SEC << endl;
+    cout << "------" << endl;
     for (int i=0; i < n; ++i)
     {
-        time_t seconds = times[2*i+1].tv_sec - times[2*i+0].tv_sec;
-        suseconds_t ms = times[2*i+1].tv_usec - times[2*i+0].tv_usec;
-        if (ms < 0) { ms += 1000000; seconds--; }
-        std::cout << labels[i] << " ::    " << (float)seconds + ms/1000000.f << std::endl;
+        time_t seconds = times[2*i+1].tv_sec  - times[2*i+0].tv_sec;
+        suseconds_t us = times[2*i+1].tv_usec - times[2*i+0].tv_usec;
+        if (us < 0) { us += 1000000; seconds--; }
+        double secs = (double)seconds + us/1000000.0f;
+        cout << labels[i] << " :: " << setiosflags(ios::fixed)<<setprecision(6)<< setw(12)<<secs<<endl;
     }
 }
 #endif
-
diff --git a/OpenMP_gpu/basic/worksharing/worksharing.f90 b/OpenMP_gpu/basic/worksharing/worksharing.f90
new file mode 100644
index 0000000..235c23b
--- /dev/null
+++ b/OpenMP_gpu/basic/worksharing/worksharing.f90
@@ -0,0 +1,33 @@
+  !! TODO 2: include  timer.f90 here
+  ...
+
+program main
+  !! TODO 2: use mod_timer here
+  use ...
+
+  implicit none
+  integer, parameter :: N=1024*1024*2**8
+  real               :: x(N), y(N), a=2.0e0; 
+  integer            :: i
+
+  !! TODO 2: Instantiate timer object, Timer_collector, as timer.
+  ...
+
+  !! initialize x and y arrays on the host
+  do i = 1,N;  x(i) = 1.0e0; y(i) = 1.0e0;  enddo
+
+  !! TODO 2: Start timer here
+  ...
+
+  !! TODO 1: Offload with teams distribute parallel for
+    !$omp  ...
+    do i = 1,N;  y(i) = a*x(i) + y(i); enddo
+
+  !! TODO 2: Stop timer here and report resuls
+  ...
+  ...
+
+end program main
+! if Intel compile error:
+!ans_worksharing.f90:(.text+0x5f): relocation truncated to fit: R_X86_64_PC32 against `.bss'
+! ifort -mcmodel=medium   *.f90