Initial support for generating an OpenCL PSy layer

doc/developers.rst

@@ -1489,3 +1489,139 @@ Of course, a given field may already be on the device (and have been
 updated) due to a previous Invoke. In this case, the fact that the
 OpenACC run-time does not copy over the now out-dated host version of
 the field is essential for correctness.
+
+.. _opencl_dev:
+
+OpenCL Support
+##############
+
+PSyclone is able to generate an OpenCL :cite:`opencl` version of
+PSy-layer code for the GOcean 1.0 API. Such code may then be executed
+on devices such as GPUs and FPGAs (Field-Programmable Gate
+Arrays). Since OpenCL code is very different to that which PSyclone
+normally generates, its creation is handled by ``gen_ocl`` methods
+instead of the normal ``gen_code``. Which of these to use is
+determined by the value of the ``Schedule.opencl`` flag.  In turn,
+this is set at a user level by the ``transformations.OCLTrans``
+transformation.
+
+The PSyKAl model of calling kernels for pre-determined iteration
+spaces is a natural fit to OpenCL's concept of an
+``NDRangeKernel``. However, the kernels themselves must be created or
+loaded at runtime, their arguments explicitly set and any arrays
+copied to the compute device. All of this 'boilerplate' code is
+generated by PSyclone. In order to minimise the changes required, the
+generated code is still Fortran and makes use of the FortCL library
+(https://github.com/stfc/FortCL) to access OpenCL functionality. We
+could of course generate the PSy layer in C instead but this would
+require further extension of PSyclone.
+
+Consider the following invoke::
+
+    call invoke( compute_cu(CU_fld, p_fld, u_fld) )
+
+When creating the OpenCL PSy layer for this invoke, PSyclone creates
+three subroutines instead of the usual one. The first, ``psy_init``
+is responsible for ensuring that a valid kernel object is created
+for each kernel called by the invoke, e.g.::
+
+    use fortcl, only: ocl_env_init, add_kernels
+    ...
+    ! Initialise the OpenCL environment/device
+    CALL ocl_env_init
+    ! The kernels this PSy layer module requires
+    kernel_names(1) = "compute_cu_code"
+    ! Create the OpenCL kernel objects. Expects to find all of the
+    ! compiled kernels in PSYCLONE_KERNELS_FILE.
+    CALL add_kernels(1, kernel_names)
+
+As indicated in the comment, the ``FortCL::add_kernels`` routine
+expects to find all kernels in a pre-compiled file pointed to by the
+PSYCLONE_KERNELS_FILE environment variable. (A pre-compiled file is
+used instead of run-time kernel compilation in order to support
+execution on FPGAs.)
+
+The second routine created by PSyclone sets the kernel arguments, e.g.::
+
+    SUBROUTINE compute_cu_code_set_args(kernel_obj, nx, cu_fld, p_fld, u_fld)
+      USE clfortran, ONLY: clSetKernelArg
+      USE iso_c_binding, ONLY: c_sizeof, c_loc, c_intptr_t
+      ...
+      INTEGER(KIND=c_intptr_t), target :: cu_fld, p_fld, u_fld
+      INTEGER(KIND=c_intptr_t), target :: kernel_obj
+      INTEGER, target :: nx
+      ! Set the arguments for the compute_cu_code OpenCL Kernel
+      ierr = clSetKernelArg(kernel_obj, 0, C_SIZEOF(nx), C_LOC(nx))
+      ierr = clSetKernelArg(kernel_obj, 1, C_SIZEOF(cu_fld), C_LOC(cu_fld))
+      ...
+    END SUBROUTINE compute_cu_code_set_args
+
+The third routine generated is the ususal psy-layer routine that is
+responsible for calling all of the kernels. However, it must now also
+call ``psy_init``, create buffers on the compute device (if they are
+not already present) and copy data over::
+
+    SUBROUTINE invoke_compute_cu(...)
+      ...
+      IF (first_time) THEN
+        first_time = .false.
+        CALL psy_init
+        num_cmd_queues = get_num_cmd_queues()
+        cmd_queues => get_cmd_queues()
+        kernel_compute_cu_code = get_kernel_by_name("compute_cu_code")
+      END IF 
+      globalsize = (/p_fld%grid%nx, p_fld%grid%ny/)
+      ! Ensure field data is on device
+      IF (.NOT. cu_fld%data_on_device) THEN
+        size_in_bytes = int(p_fld%grid%nx*p_fld%grid%ny, 8)* &
+                        c_sizeof(cu_fld%data(1,1))
+        ! Create buffer on device
+        cu_fld%device_ptr = create_rw_buffer(size_in_bytes)
+        ierr = clEnqueueWriteBuffer(cmd_queues(1), cu_fld%device_ptr,  &
+                                    CL_TRUE, 0_8, size_in_bytes,       &
+      			            C_LOC(cu_fld%data), 0, C_NULL_PTR, &
+      			            C_LOC(write_event))
+        cu_fld%data_on_device = .true.
+      END IF 
+      ...
+
+Note that we use the ``data_on_device`` member of the field derived
+type (implemented in github.com/stfc/dl_esm_inf) to keep track of
+whether a given field has been copied to the compute device.  Once all
+of this setup is done, the kernel itself is launched by calling
+``clEnqueueNDRangeKernel``::
+
+    ierr = clEnqueueNDRangeKernel(cmd_queues(1), kernel_compute_cu_code, &
+                                  2, C_NULL_PTR, C_LOC(globalsize),      &
+				  C_NULL_PTR, 0, C_NULL_PTR, C_NULL_PTR)
+
+Limitations
+===========
+
+Currently PSyclone can only generate the OpenCL version of the PSy
+layer.  Execution of the resulting code requires that the kernels
+themselves be converted from Fortran to OpenCL (a dialect of C) and at
+present this must be done manually. Since all data accessed by an
+OpenCL kernel must be passed as an argument, this conversion must also
+convert any accesses to module data into routine arguments. 
+Work is in progress to support kernel transformation and this will be
+made available in a future PSyclone release.
+
+In OpenCL, all tasks to be performed (whether copying data or kernel
+execution) are associated with a command queue. Tasks submitted to
+different command queues may then be executed concurrently,
+potentially giving greater performance. The OpenCL PSy code currently
+generated by PSyclone makes use of just one command queue but again,
+this could be extended in the future.
+
+The current implementation only supports the conversion of a whole
+Invoke to use OpenCL. In the future we may refine this functionality
+so that it may be applied to just a subset of kernels within an
+Invoke.
+
+Since PSyclone knows nothing about the I/O performed by a model, the
+task of ensuring that the correct data is written out by a model
+(including when doing halo exchanges for distributed memory) is left
+to the dl_esm_inf library since that has the information on whether
+field data is local or on a remote compute device.
+

doc/references.bib

@@ -14,3 +14,10 @@ @manual{nemo_code_conv
    url = "https://forge.ipsl.jussieu.fr/nemo/attachment/wiki/Literature/NEMO_coding.conv_v3.pdf",
    year = "2013"
 }
+
+@manual{opencl,
+   title = "The OpenCL 2.2 Reference Guide",
+   version = "2.2",
+   url = "https://www.khronos.org/files/opencl22-reference-guide.pdf",
+   year = "2018"
+}

doc/transformations.rst

@@ -35,9 +35,9 @@ The generic transformations currently available are listed in
 alphabetical order below (a number of these have specialisations which
 can be found in the API-specific sections).
 
-.. note:: PSyclone currently only supports OpenACC transformations
-	  for the GOcean 1.0 API. Attempts to apply these
-	  transformations to (members of) Schedules from other
+.. note:: PSyclone currently only supports OpenACC and OpenCL
+	  transformations for the GOcean 1.0 API. Attempts to apply
+	  these transformations to (members of) Schedules from other
 	  APIs will be rejected.
 
 ####
@@ -90,9 +90,12 @@ can be found in the API-specific sections).
 
 ####
 
-.. autoclass:: psyclone.transformations.ProfileRegionTrans
-    :members:
-    :noindex:
+.. autoclass:: psyclone.transformations.OCLTrans
+      :members:
+      :noindex:
+
+.. note:: OpenCL support is still under development. See
+          :ref:`opencl_dev` for more details.
 
 ####
 
@@ -122,6 +125,13 @@ can be found in the API-specific sections).
 	  :ref:`MoveTrans <sec_move_trans>` transformation may be used
           for this.
 
+####
+
+.. autoclass:: psyclone.transformations.ProfileRegionTrans
+    :members:
+    :noindex:
+
+
 Kernels
 -------
 
@@ -386,3 +396,29 @@ region for a set of nodes that includes halo swaps or global sums will
 produce an error.  In such cases it may be possible to re-order the
 nodes in the Schedule using the :ref:`MoveTrans <sec_move_trans>`
 transformation.
+
+OpenCL
+------
+
+In common with OpenMP, the conversion of the generated code to use
+OpenCL is performed by a transformation (``OCLTrans`` - see the
+:ref:`sec_transformations_available` Section above). Currently this
+transformation is only supported for the GOcean1.0 API and is applied
+to the whole Schedule of an Invoke. This means that all kernels in
+that Invoke will be executed on the OpenCL device. At present the
+``OCLTrans`` transformation only alters the generated PSy-layer code. It
+is currently the user's responsibility to convert the actual kernel code
+from Fortran into OpenCL. Work is underway to extend PSyclone in
+order to perform this translation automatically.
+
+The OpenCL code generated by PSyclone is still Fortran and makes use
+of the FortCL library (https://github.com/stfc/FortCL) to access
+OpenCL functionality.  It also relies upon the OpenCL support provided
+by the dl_esm_inf library (https://github.com/stfc/dl_esm_inf).
+
+The introduction of OpenCL code generation in PSyclone has been
+largely motivated by the need to target Field Programmable Gate Array
+(FPGA) accelerator devices. It is not currently designed to target the other
+compute devices that OpenCL supports (such as GPUs and multi-core CPUs) but
+this is a potentially fruitful area for future work.
+

examples/gocean/eg3/README

@@ -0,0 +1,59 @@
+# -----------------------------------------------------------------------------
+# BSD 3-Clause License
+#
+# Copyright (c) 2018, Science and Technology Facilities Council
+# All rights reserved.
+#
+# Redistribution and use in source and binary forms, with or without
+# modification, are permitted provided that the following conditions are met:
+#
+# * Redistributions of source code must retain the above copyright notice, this
+#   list of conditions and the following disclaimer.
+#
+# * Redistributions in binary form must reproduce the above copyright notice,
+#   this list of conditions and the following disclaimer in the documentation
+#   and/or other materials provided with the distribution.
+#
+# * Neither the name of the copyright holder nor the names of its
+#   contributors may be used to endorse or promote products derived from
+#   this software without specific prior written permission.
+#
+# THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS
+# "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT
+# LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS
+# FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE
+# COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT,
+# INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING,
+# BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+# LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+# CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT
+# LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN
+# ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE
+# POSSIBILITY OF SUCH DAMAGE.
+#------------------------------------------------------------------------------
+# Author A. R. Porter, STFC Daresbury Lab
+
+The directory containing this file contains an example of the use of
+PSyclone to generate OpenCL driver code with the GOcean 1.0 API.
+
+In order to use PSyclone you must first install it, ideally with pip.
+See ../../../README.md for more details.
+
+PSyclone can be run in the directory containing this file by 
+executing, e.g.::
+
+    psyclone -api "gocean1.0" alg.f90
+
+This will generate 'vanilla' PSy-layer code which is output to stdout.
+
+In order to generate an OpenCL PSy layer instead, PSyclone must be
+provided with a transformation script::
+
+    psyclone -api "gocean1.0" -s ./ocl_trans.py alg.f90
+
+where ocl_trans.py simply applies the psyclone.transformations.OCLTrans
+transformation to the Schedule of the Invoke.
+
+Currently the (Fortran) kernels called by the Invoke must be manually
+translated into OpenCL. This step will be automated in a future
+release of PSyclone.

examples/gocean/eg3/alg.f90

@@ -0,0 +1,78 @@
+! -----------------------------------------------------------------------------
+! BSD 3-Clause License
+!
+! Copyright (c) 2018, Science and Technology Facilities Council.
+! All rights reserved.
+!
+! Redistribution and use in source and binary forms, with or without
+! modification, are permitted provided that the following conditions are met:
+!
+! * Redistributions of source code must retain the above copyright notice, this
+!   list of conditions and the following disclaimer.
+!
+! * Redistributions in binary form must reproduce the above copyright notice,
+!   this list of conditions and the following disclaimer in the documentation
+!   and/or other materials provided with the distribution.
+!
+! * Neither the name of the copyright holder nor the names of its
+!   contributors may be used to endorse or promote products derived from
+!   this software without specific prior written permission.
+!
+! THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+! AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+! IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+! DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+! FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+! DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+! SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+! CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+! OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+! OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+! -----------------------------------------------------------------------------
+! Author: A. R. Porter, STFC Daresbury Lab.
+
+! A simple, single Invoke example to demonstrate the generation of an
+! OpenCL driver PSy layer.
+program simple
+
+  use grid_mod
+  use field_mod
+  use compute_cu_mod,   only: compute_cu
+  use compute_cv_mod,   only: compute_cv
+  use compute_z_mod,    only: compute_z
+  use compute_h_mod,    only: compute_h
+  implicit none
+
+  type(grid_type), target :: model_grid
+
+  type(r2d_field) :: p_fld
+  type(r2d_field) :: u_fld, v_fld
+  type(r2d_field) :: cu_fld, cv_fld
+  type(r2d_field) :: z_fld
+  type(r2d_field) :: h_fld
+
+  integer :: ncycle
+
+  model_grid = grid_type(ARAKAWA_C,                           &
+                         (/BC_PERIODIC,BC_PERIODIC,BC_NONE/), &
+                         OFFSET_SW)
+
+  ! Create fields on this grid
+  p_fld  = r2d_field(model_grid, T_POINTS)
+  u_fld  = r2d_field(model_grid, U_POINTS)
+  v_fld  = r2d_field(model_grid, V_POINTS)
+  cu_fld = r2d_field(model_grid, U_POINTS)
+  cv_fld = r2d_field(model_grid, V_POINTS)
+  z_fld = r2d_field(model_grid, F_POINTS)
+  h_fld = r2d_field(model_grid, T_POINTS)
+
+  do ncycle=1,itmax
+
+    call invoke( compute_cu(CU_fld, p_fld, u_fld),      &
+                 compute_cv(CV_fld, p_fld, v_fld),      &
+                 compute_z(z_fld, p_fld, u_fld, v_fld), &
+                 compute_h(h_fld, p_fld, u_fld, v_fld) )
+
+  end do
+
+end program simple