diff --git a/History.tex b/History.tex index a1eabb1..8de34d9 100644 --- a/History.tex +++ b/History.tex @@ -18,7 +18,7 @@ \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 @@ \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} diff --git a/devices/declare_target.tex b/devices/declare_target.tex index febbccc..c83d352 100644 --- a/devices/declare_target.tex +++ b/devices/declare_target.tex @@ -246,3 +246,25 @@ double precision data are in scope on the device only for the second \code{targe \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} + diff --git a/devices/sources/declare_target.7.c b/devices/sources/declare_target.7.c new file mode 100644 index 0000000..4e9d9ea --- /dev/null +++ b/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 + +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; + +} diff --git a/devices/sources/declare_target.7.f90 b/devices/sources/declare_target.7.f90 new file mode 100644 index 0000000..ff9eccd --- /dev/null +++ b/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 diff --git a/tasking/sources/tasking.8.f b/tasking/sources/tasking.8.f index 30b15d6..cbaa60d 100644 --- a/tasking/sources/tasking.8.f +++ b/tasking/sources/tasking.8.f @@ -7,8 +7,10 @@ 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 @@ !$omp end task !$omp end parallel end subroutine - end module