declare target with device_type(nohost)

Author: Henry Jin

History.tex

@@ -18,7 +18,7 @@ \section{Changes from 5.2 to 5.2.1}
 \item Added the following examples for the 5.2 features:
 \begin{itemize}
   \item \scode{uses_allocators} clause for the use of allocators in 
-  \code{target} regions (\specref{sec:allocators})
+     \code{target} regions (\specref{sec:allocators})
 \end{itemize}
 \item Added the following examples for the 5.1 features:
 \begin{itemize}
@@ -27,6 +27,8 @@ \section{Changes from 5.2 to 5.2.1}
 \end{itemize}
 \item Added other examples:
 \begin{itemize}
+  \item \code{declare}~\code{target} directive with \scode{device_type(nohost)}
+     clause (\specref{subsec:declare_target_device_type})
   \item \scode{omp_pause_resource} and \scode{omp_pause_resource_all}
      routines (\specref{sec:pause_resource})
 \end{itemize}

devices/declare_target.tex

@@ -246,3 +246,25 @@ \subsection{Declare Target Directive with \code{link} Clause}
 \cexample[5.1]{declare_target}{6}
 \ffreeexample[4.5]{declare_target}{6}
 
+
+\subsection{Declare Target Directive with \code{device\_type} Clause}
+\label{subsec:declare_target_device_type}
+
+\index{directives!declare target@\code{declare}~\code{target}}
+\index{declare target directive@\code{declare}~\code{target} directive}
+
+\index{directives!begin declare target@\code{begin}~\code{declare}~\code{target}}
+\index{begin declare target directive@\code{begin}~\code{declare}~\code{target} directive}
+
+\index{clauses!device_type@\code{device\_type}}
+\index{device_type clause@\code{device\_type} clause}
+
+The \code{declare}~\code{target} directives apply to procedures to ensure that they can be executed or accessed on a device. 
+The \code{device\_type} clause specifies whether a version of the procedure or variable should be made available on the host, device or both. 
+This example uses  \code{nohost} for a procedure \plc{foo}. Only a device version of the procedure \plc{foo} is made available. 
+If the variant function \plc{foo\_onhost} is not specified for the host fallback execution, the call to \plc{foo} from the \code{target} region will result in a link time error due to the code generated for host execution of the target region.
+This is because host symbol for the device routine \plc{foo} marked as \code{nohost} is not required to be present in the host environment.
+
+\cexample[5.2]{declare_target}{7}
+\ffreeexample[5.2]{declare_target}{7}
+

devices/sources/declare_target.7.c

@@ -0,0 +1,32 @@
+/*
+* @@name: declare_target.7
+* @@type: C
+* @@operation: link
+* @@expect: success
+* @@version: omp_5.2
+*/
+#include <stdio.h>
+
+void foo();
+void foo_onhost();
+
+#pragma omp declare target enter(foo) device_type(nohost)
+
+#pragma omp declare variant(foo_onhost) match(device={kind(host)})
+void foo(){
+  // device specific computation
+}
+
+void foo_onhost(){
+  printf("On host\n");
+} 
+
+int main(){
+  #pragma omp target teams 
+  {
+    foo();  // calls foo() on target device 
+            // or foo_onhost() in case of host fallback
+  }
+  return 0;
+
+}

devices/sources/declare_target.7.f90

@@ -0,0 +1,29 @@
+! @@name: declare_target.7
+! @@type: F-free
+! @@operation: link 
+! @@expect: success
+! @@version: omp_5.2
+module subs
+
+contains
+   subroutine foo()
+   !$omp declare target enter(foo) device_type(nohost)
+   !$omp declare variant(foo_onhost) match(device={kind(host)})
+     ! device specific computation
+   end subroutine
+
+   subroutine foo_onhost()
+     print *,' On host.'
+   end subroutine
+
+end module
+
+program main
+
+   use subs
+   !$omp target
+   call foo    ! calls foo() on device
+               ! or foo_onhost() in case of host fallback
+   !$omp end target
+
+end program

tasking/sources/tasking.8.f

@@ -7,8 +7,10 @@ module example
       integer tp
 !$omp threadprivate(tp)
       integer var
-      contains
+      end module
+
       subroutine work
+         use example
 !$omp parallel
          ! do work here
 !$omp task
@@ -21,4 +23,3 @@ subroutine work
 !$omp end task
 !$omp end parallel
       end subroutine
-      end module