Inheritance diagram for psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans:

[legend]

Collaboration diagram for psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans:

[legend]

Public Member Functions
def	name (self)

def	validate (self, node, options=None)

def	apply (self, node, options=None)

Detailed Description

Switches on/off the generation of an OpenCL PSy layer for a given
InvokeSchedule. Additionally, it will generate OpenCL kernels for
each of the kernels referenced by the Invoke. For example:

>>> from psyclone.parse.algorithm import parse
>>> from psyclone.psyGen import PSyFactory
>>> API = "gocean1.0"
>>> FILENAME = "shallow_alg.f90" # examples/gocean/eg1
>>> ast, invoke_info = parse(FILENAME, api=API)
>>> psy = PSyFactory(API, distributed_memory=False).create(invoke_info)
>>> schedule = psy.invokes.get('invoke_0').schedule
>>> ocl_trans = GOOpenCLTrans()
>>> ocl_trans.apply(schedule)
>>> print(schedule.view())

Definition at line 60 of file gocean_opencl_trans.py.

Member Function Documentation

◆ apply()

def psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans.apply	(	self,
		node,
		options = `None`
	)

Apply the OpenCL transformation to the supplied GOInvokeSchedule. This
causes PSyclone to generate an OpenCL version of the corresponding
PSy-layer routine. The generated code makes use of the FortCL
library (https://github.com/stfc/FortCL) in order to manage the
OpenCL device directly from Fortran.

:param node: the InvokeSchedule to transform.
:type node: :py:class:`psyclone.psyGen.GOInvokeSchedule`
:param options: set of option to tune the OpenCL generation.
:type options: dict of str:value or None
:param bool options["enable_profiling"]: whether or not to set up the \
        OpenCL environment with the profiling option enabled.
:param bool options["out_of_order"]: whether or not to set up the \
        OpenCL environment with the out_of_order option enabled.
:param bool options["end_barrier"]: whether or not to add an OpenCL \
        barrier at the end of the transformed invoke.

Reimplemented from psyclone.psyGen.Transformation.

Definition at line 222 of file gocean_opencl_trans.py.

     def apply(self, node, options=None):
         '''
         Apply the OpenCL transformation to the supplied GOInvokeSchedule. This
         causes PSyclone to generate an OpenCL version of the corresponding
         PSy-layer routine. The generated code makes use of the FortCL
         library (https://github.com/stfc/FortCL) in order to manage the
         OpenCL device directly from Fortran.
  
         :param node: the InvokeSchedule to transform.
         :type node: :py:class:`psyclone.psyGen.GOInvokeSchedule`
         :param options: set of option to tune the OpenCL generation.
         :type options: dict of str:value or None
         :param bool options["enable_profiling"]: whether or not to set up the \
                 OpenCL environment with the profiling option enabled.
         :param bool options["out_of_order"]: whether or not to set up the \
                 OpenCL environment with the out_of_order option enabled.
         :param bool options["end_barrier"]: whether or not to add an OpenCL \
                 barrier at the end of the transformed invoke.
  
         '''
         if not options:
             options = {}
  
         self.validate(node, options)
         api_config = Config.get().api_conf("gocean1.0")
  
         # Update class attributes
         if 'enable_profiling' in options:
             self._enable_profiling = options['enable_profiling']
  
         if 'out_of_order' in options:
             self._out_of_order = options['out_of_order']
  
         self._transformed_invokes += 1
  
         # Get end_barrier option
         end_barrier = options.get('end_barrier', True)
  
         # Update the maximum value that the queue_number have.
         for kernel in node.coded_kernels():
             self._max_queue_number = max(self._max_queue_number,
                                          kernel.opencl_options["queue_number"])
  
         # Insert, if they don't already exist, the necessary OpenCL helper
         # subroutines in the root Container.
         psy_init = self._insert_opencl_init_routine(node.root)
         init_grid = self._insert_initialise_grid_buffers(node.root)
         write_grid_buf = self._insert_write_grid_buffers(node.root)
         self._insert_ocl_read_from_device_function(node.root)
         self._insert_ocl_write_to_device_function(node.root)
         init_buf = self._insert_ocl_initialise_buffer(node.root)
  
         for kern in node.coded_kernels():
             self._insert_ocl_arg_setter_routine(node.root, kern)
  
         # Insert fortcl, clfortran and c_iso_binding import statement
         fortcl = ContainerSymbol("fortcl")
         node.symbol_table.add(fortcl)
         get_num_cmd_queues = RoutineSymbol(
                 "get_num_cmd_queues", interface=ImportInterface(fortcl))
         get_cmd_queues = RoutineSymbol(
                 "get_cmd_queues", interface=ImportInterface(fortcl))
         get_kernel_by_name = RoutineSymbol(
                 "get_kernel_by_name", interface=ImportInterface(fortcl))
         node.symbol_table.add(get_num_cmd_queues)
         node.symbol_table.add(get_cmd_queues)
         node.symbol_table.add(get_kernel_by_name)
         clfortran = ContainerSymbol("clfortran")
         node.symbol_table.add(clfortran)
         cl_finish = RoutineSymbol(
                 "clFinish", interface=ImportInterface(clfortran))
         cl_launch = RoutineSymbol(
                 "clEnqueueNDRangeKernel",
                 interface=ImportInterface(clfortran))
         node.symbol_table.add(cl_finish)
         node.symbol_table.add(cl_launch)
         iso_c_binding = ContainerSymbol("iso_c_binding")
         node.symbol_table.add(iso_c_binding)
         c_loc = RoutineSymbol(
                 "C_LOC", interface=ImportInterface(iso_c_binding))
         c_null = DataSymbol(
                 "C_NULL_PTR", datatype=INTEGER_TYPE,
                 interface=ImportInterface(iso_c_binding))
         node.symbol_table.add(c_loc)
         node.symbol_table.add(c_null)
  
         # Include the check_status subroutine if we are in debug_mode
         if api_config.debug_mode:
             ocl_utils = ContainerSymbol("ocl_utils_mod")
             check_status = RoutineSymbol(
                 "check_status", interface=ImportInterface(ocl_utils))
             node.symbol_table.add(ocl_utils)
             node.symbol_table.add(check_status)
  
         # Declare local variables needed by an OpenCL PSy-layer invoke
         qlist = node.symbol_table.new_symbol(
             "cmd_queues", symbol_type=DataSymbol,
             datatype=UnsupportedFortranType(
                 "integer(kind=c_intptr_t), pointer, save :: cmd_queues(:)"),
             tag="opencl_cmd_queues")
         # 'first_time' needs to be an UnsupportedFortranType because it has
         # SAVE and initial value
         first = DataSymbol("first_time",
                            datatype=UnsupportedFortranType(
                                "logical, save :: first_time = .true."))
         node.symbol_table.add(first, tag="first_time")
         flag = node.symbol_table.new_symbol(
             "ierr", symbol_type=DataSymbol, datatype=INTEGER_TYPE,
             tag="opencl_error")
         global_size = node.symbol_table.new_symbol(
             "globalsize", symbol_type=DataSymbol,
             datatype=UnsupportedFortranType(
                 "integer(kind=c_size_t), target :: globalsize(2)"))
         local_size = node.symbol_table.new_symbol(
             "localsize", symbol_type=DataSymbol,
             datatype=UnsupportedFortranType(
                 "integer(kind=c_size_t), target :: localsize(2)"))
  
         # Bring all the boundaries at the beginning (since we are going to
         # use them during the setup block - and they don't change)
         boundary_vars = []
         for tag, symbol in node.symbol_table.tags_dict.items():
             if tag.startswith(("xstart_", "xstop_", "ystart_", "ystop_")):
                 boundary_vars.append(symbol)
         cursor = 0
         for assignment in node.walk(Assignment):
             if assignment.lhs.symbol in boundary_vars:
                 node.children.insert(cursor, assignment.detach())
                 cursor += 1
  
         # Create block of code to execute only the first time:
         setup_block = IfBlock.create(Reference(first), [])
         setup_block.preceding_comment = \
             "Initialise OpenCL runtime, kernels and buffers"
         node.children.insert(cursor, setup_block)
         setup_block.if_body.addchild(Call.create(psy_init, []))
  
         # Set up cmd_queues pointer
         ptree = Fortran2003.Pointer_Assignment_Stmt(
             f"{qlist.name} => {get_cmd_queues.name}()")
         cblock = CodeBlock([ptree], CodeBlock.Structure.STATEMENT)
         setup_block.if_body.addchild(cblock)
  
         # Declare and assign kernel pointers
         for kern in node.coded_kernels():
             name = "kernel_" + kern.name
             try:
                 kpointer = node.symbol_table.lookup_with_tag(name)
             except KeyError:
                 pointer_type = UnsupportedFortranType(
                     "INTEGER(KIND=c_intptr_t), TARGET, SAVE :: " + name)
                 kpointer = DataSymbol(name, datatype=pointer_type)
                 node.symbol_table.add(kpointer, tag=name)
             setup_block.if_body.addchild(
                 Assignment.create(
                     Reference(kpointer),
                     Call.create(get_kernel_by_name,
                                 [Literal(kern.name, CHARACTER_TYPE)])))
  
         # Traverse all arguments and make sure all the buffers are initialised
         initialised_fields = set()
         there_is_a_grid_buffer = False
         for kern in node.coded_kernels():
             for arg in kern.arguments.args:
                 if arg.argument_type == "field":
                     field = node.symbol_table.lookup(arg.name)
                     if field not in initialised_fields:
                         # Call the init_buffer routine with this field
                         call = Call.create(init_buf, [Reference(field)])
                         setup_block.if_body.addchild(call)
                         initialised_fields.add(field)
                 elif (arg.argument_type == "grid_property" and
                       not arg.is_scalar):
                     if not there_is_a_grid_buffer:
                         # Call the grid init_buffer routine
                         field = node.symbol_table.lookup(
                                 kern.arguments.find_grid_access().name)
                         call = Call.create(init_grid, [Reference(field)])
                         setup_block.if_body.addchild(call)
                         there_is_a_grid_buffer = True
                 if not arg.is_scalar:
                     # All buffers will be assigned to a local OpenCL memory
                     # object to easily reference them, make sure this local
                     # variable is declared in the Invoke.
                     name = arg.name + "_cl_mem"
                     try:
                         node.symbol_table.lookup_with_tag(name)
                     except KeyError:
                         node.symbol_table.new_symbol(
                             name, tag=name, symbol_type=DataSymbol,
                             datatype=UnsupportedFortranType(
                                 "INTEGER(KIND=c_intptr_t) :: " + name))
  
         # Now call all the set_args routines because in some platforms (e.g.
         # in Xilinx FPGA) knowing which arguments each kernel is going to use
         # allows the write operation to place the data into the appropriate
         # memory bank.
         first_statement_comment = False
         kernel_names = set()
         for kern in node.coded_kernels():
             if kern.name not in kernel_names:
                 kernel_names.add(kern.name)
                 callblock = self._generate_set_args_call(kern, node.scope)
                 for child in callblock.pop_all_children():
                     setup_block.if_body.addchild(child)
                     if not first_statement_comment:
                         child.preceding_comment = (
                             "Do a set_args now so subsequent writes place the "
                             "data appropriately")
                         first_statement_comment = True
  
         # Now we can insert calls to write_to_device method for each buffer
         # and the grid writing call if there is one (in a new first time block)
         first_statement_comment = False
         for field in initialised_fields:
             call = Call.create(
                 RoutineSymbol(field.name+"%write_to_device"), [])
             setup_block.if_body.addchild(call)
             if not first_statement_comment:
                 call.preceding_comment = "Write data to the device"
                 first_statement_comment = True
  
         if there_is_a_grid_buffer:
             fieldarg = node.coded_kernels()[0].arguments.find_grid_access()
             field = node.symbol_table.lookup(fieldarg.name)
             call = Call.create(write_grid_buf, [Reference(field)])
             setup_block.if_body.addchild(call)
  
         # We will just mark the nodes we are replacing as deleting them inside
         # the loop would break the PSy-layer backward_dependency method in the
         # following iterations. We will detach all these nodes after the loop.
         nodes_to_detach = []
  
         # Transform each kernel call loop construct to its equivalent FortCL
         # statements
         for kern in node.coded_kernels():
             outerloop = kern.ancestor(GOLoop).ancestor(GOLoop)
  
             # Set up globalsize and localsize arrays
             garg = node.coded_kernels()[0].arguments.find_grid_access()
             num_x = api_config.grid_properties["go_grid_nx"].fortran\
                 .format(garg.name)
             num_y = api_config.grid_properties["go_grid_ny"].fortran\
                 .format(garg.name)
             assig = Assignment.create(
                     Reference(global_size),
                     Literal(f"(/{num_x}, {num_y}/)",
                             ArrayType(INTEGER_TYPE, [2])))
             node.children.insert(outerloop.position, assig)
             local_size_value = kern.opencl_options['local_size']
             assig = Assignment.create(
                     Reference(local_size),
                     Literal(f"(/{local_size_value}, 1/)",
                             ArrayType(INTEGER_TYPE, [2])))
             node.children.insert(outerloop.position, assig)
  
             # Check that the global_size is multiple of the local_size
             if api_config.debug_mode:
                 fortran_reader = FortranReader()
                 global_size_expr = fortran_reader.psyir_from_expression(
                         num_x, node.symbol_table)
                 self._add_divisibility_check(node, outerloop.position,
                                              check_status, global_size_expr,
                                              local_size_value)
  
             # Retrieve kernel symbol
             kernelsym = node.symbol_table.lookup_with_tag(
                             "kernel_" + kern.name)
  
             # Choose the command queue number to which to dispatch this kernel.
             # We have do deal with possible dependencies to kernels dispatched
             # in different command queues as the order of execution is not
             # guaranteed.
             queue_number = kern.opencl_options['queue_number']
             cmd_queue = ArrayReference.create(
                     qlist, [Literal(str(queue_number), INTEGER_TYPE)])
             dependency = outerloop.backward_dependence()
  
             # If the dependency is a loop containing a kernel, add a barrier if
             # the previous kernels were dispatched in a different command queue
             if dependency:
                 for kernel_dep in dependency.coded_kernels():
                     previous_queue = kernel_dep.opencl_options['queue_number']
                     if previous_queue != queue_number:
                         # If the backward dependency is being executed in
                         # another queue we add a barrier to make sure the
                         # previous kernel has finished before this halo
                         # exchange starts.
                         barrier = Assignment.create(
                                     Reference(flag),
                                     Call.create(cl_finish, [
                                         ArrayReference.create(qlist, [
                                             Literal(str(previous_queue),
                                                     INTEGER_TYPE)])]))
                         node.children.insert(outerloop.position, barrier)
  
             # If the dependency is something other than a kernel, currently we
             # dispatch everything else to queue _OCL_MANAGEMENT_QUEUE, so add a
             # barrier if this kernel is not on queue _OCL_MANAGEMENT_QUEUE.
             if dependency and not dependency.coded_kernels() and \
                     queue_number != self._OCL_MANAGEMENT_QUEUE:
                 barrier = Assignment.create(
                             Reference(flag),
                             Call.create(cl_finish, [
                                 ArrayReference.create(qlist, [
                                     Literal(str(self._OCL_MANAGEMENT_QUEUE),
                                             INTEGER_TYPE)])]))
                 node.children.insert(outerloop.position, barrier)
  
             # Check that everything has succeeded before the kernel launch
             if api_config.debug_mode:
                 self._add_ready_check(node, outerloop.position, check_status,
                                       kern.name, flag, cl_finish,
                                       cmd_queue.copy())
             callblock = self._generate_set_args_call(kern, node.scope)
             for child in callblock.pop_all_children():
                 node.children.insert(outerloop.position, child)
  
             # Then we call the clEnqueueNDRangeKernel
             assig = Assignment.create(
                         Reference(flag),
                         Call.create(cl_launch, [
                             # OpenCL Command Queue
                             cmd_queue,
                             # OpenCL Kernel object
                             Reference(kernelsym),
                             # Number of work dimensions
                             Literal("2", INTEGER_TYPE),
                             # Global offset (if NULL the global IDs start at
                             # offset (0,0,0))
                             Reference(c_null),
                             # Global work size
                             Call.create(c_loc, [Reference(global_size)]),
                             # Local work size
                             Call.create(c_loc, [Reference(local_size)]),
                             # Number of events in wait list
                             Literal("0", INTEGER_TYPE),
                             # Event wait list that need to be completed before
                             # this kernel
                             Reference(c_null),
                             # Event that identifies this kernel completion
                             Reference(c_null)]))
             assig.preceding_comment = "Launch the kernel"
             node.children.insert(outerloop.position, assig)
             self._insert_kernel_code_in_opencl_file(kern)
  
             # Add additional checks if we are in debug mode
             if api_config.debug_mode:
                 self._add_kernel_check(node, outerloop.position, check_status,
                                        kern.name, flag, cl_finish,
                                        cmd_queue.copy())
  
             nodes_to_detach.append(outerloop)
  
         # If we execute the kernels asynchronously, we need to add wait
         # statements before the halo exchanges to guarantee that the data
         # has been updated
         for possible_dependent_node in node.walk(HaloExchange):
             # The backward_dependences returns the last Loop with a kernel
             # that has a dependency with this halo exchange
             dependency = possible_dependent_node.backward_dependence()
             if dependency:
                 for kernel_dep in dependency.coded_kernels():
                     previous_queue = kernel_dep.opencl_options['queue_number']
                     if previous_queue != self._OCL_MANAGEMENT_QUEUE:
                         # If the backward dependency is being executed in
                         # another queue we add a barrier to make sure the
                         # previous kernel has finished before this one starts.
                         barrier = Assignment.create(
                                     Reference(flag),
                                     Call.create(cl_finish, [
                                         ArrayReference.create(qlist, [
                                             Literal(str(previous_queue),
                                                     INTEGER_TYPE)])]))
                         pos = possible_dependent_node.position
                         node.children.insert(pos, barrier)
  
         for node_to_detach in nodes_to_detach:
             node_to_detach.detach()
  
         if end_barrier:
             self._add_end_barrier(node, flag, cl_finish, qlist)
  
         # And at the very end always makes sure that first_time value is False
         assign = Assignment.create(Reference(first),
                                    Literal("false", BOOLEAN_TYPE))
         assign.preceding_comment = "Unset the first time flag"
         node.addchild(assign)
  
         self._output_opencl_kernels_file()
  

◆ name()

def psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans.name ( self )

:returns: the name of this transformation.
:rtype: str

Reimplemented from psyclone.psyGen.Transformation.

Definition at line 100 of file gocean_opencl_trans.py.

     def name(self):
         '''
         :returns: the name of this transformation.
         :rtype: str
         '''
         return "GOOpenCLTrans"
  

Here is the caller graph for this function:

◆ validate()

def psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans.validate	(	self,
		node,
		options = `None`
	)

Checks that the supplied InvokeSchedule is valid and that an OpenCL
version of it can be generated.

:param node: the Schedule to check.
:type node: :py:class:`psyclone.psyGen.InvokeSchedule`
:param options: a dictionary with options for transformations.
:type options: dict of str:value or None
:param bool options["enable_profiling"]: whether or not to set up the \
        OpenCL environment with the profiling option enabled.
:param bool options["out_of_order"]: whether or not to set up the \
        OpenCL environment with the out_of_order option enabled.
:param bool options["end_barrier"]: whether or not to add an OpenCL \
        barrier at the end of the transformed invoke.

:raises TransformationError: if the InvokeSchedule is not for the \
                             GOcean1.0 API.
:raises TransformationError: if any of the kernels have arguments \
                             which are passed as a literal.
:raises TransformationError: if any of the provided options is invalid.
:raises TransformationError: if any of the provided options is not \
                             compatible with a previous OpenCL
                             environment.
:raises TransformationError: if any kernel in this invoke has a \
                             global variable used by an import.
:raises TransformationError: if any kernel does not iterate over \
                             the whole grid.

Reimplemented from psyclone.psyGen.Transformation.

Definition at line 107 of file gocean_opencl_trans.py.

     def validate(self, node, options=None):
         '''
         Checks that the supplied InvokeSchedule is valid and that an OpenCL
         version of it can be generated.
  
         :param node: the Schedule to check.
         :type node: :py:class:`psyclone.psyGen.InvokeSchedule`
         :param options: a dictionary with options for transformations.
         :type options: dict of str:value or None
         :param bool options["enable_profiling"]: whether or not to set up the \
                 OpenCL environment with the profiling option enabled.
         :param bool options["out_of_order"]: whether or not to set up the \
                 OpenCL environment with the out_of_order option enabled.
         :param bool options["end_barrier"]: whether or not to add an OpenCL \
                 barrier at the end of the transformed invoke.
  
         :raises TransformationError: if the InvokeSchedule is not for the \
                                      GOcean1.0 API.
         :raises TransformationError: if any of the kernels have arguments \
                                      which are passed as a literal.
         :raises TransformationError: if any of the provided options is invalid.
         :raises TransformationError: if any of the provided options is not \
                                      compatible with a previous OpenCL
                                      environment.
         :raises TransformationError: if any kernel in this invoke has a \
                                      global variable used by an import.
         :raises TransformationError: if any kernel does not iterate over \
                                      the whole grid.
         '''
  
         if isinstance(node, InvokeSchedule):
             if not isinstance(node, GOInvokeSchedule):
                 raise TransformationError(
                     f"OpenCL generation is currently only supported for the "
                     f"GOcean API but got an InvokeSchedule of type: "
                     f"'{type(node).__name__}'")
         else:
             raise TransformationError(
                 f"Error in GOOpenCLTrans: the supplied node must be a (sub-"
                 f"class of) InvokeSchedule but got {type(node)}")
  
         # Validate options map
         valid_options = ['end_barrier', 'enable_profiling', 'out_of_order']
         for key, value in options.items():
             if key in valid_options:
                 # All current options should contain boolean values
                 if not isinstance(value, bool):
                     raise TransformationError(
                         f"InvokeSchedule OpenCL option '{key}' should be a "
                         f"boolean.")
             else:
                 raise TransformationError(
                     f"InvokeSchedule does not support the OpenCL option "
                     f"'{key}'. The supported options are: {valid_options}.")
  
         # Validate that the options are valid with previously generated OpenCL
         if self._transformed_invokes > 0:
             if ('enable_profiling' in options and
                     self._enable_profiling != options['enable_profiling']):
                 raise TransformationError(
                     f"Can't generate an OpenCL Invoke with enable_profiling='"
                     f"{options['enable_profiling']}' since a previous "
                     f"transformation used a different value, and their OpenCL"
                     f" environments must match.")
  
             if ('out_of_order' in options and
                     self._out_of_order != options['out_of_order']):
                 raise TransformationError(
                     f"Can't generate an OpenCL Invoke with out_of_order='"
                     f"{options['out_of_order']}' since a previous "
                     f"transformation used a different value, and their OpenCL "
                     f"environments must match.")
  
         # Now we need to check that none of the invoke arguments is a literal
         args = args_filter(node.args, arg_types=["scalar"])
         for arg in args:
             if arg.is_literal:
                 raise TransformationError(
                     f"Cannot generate OpenCL for Invokes that contain kernel "
                     f"arguments which are a literal, but found the literal "
                     f"'{arg.name}' used as an argument in invoke "
                     f"'{node.name}'.")
  
         # Check that we can construct the PSyIR and SymbolTable of each of
         # the kernels in this Schedule. Also check that none of them access
         # any form of global data (that is not a routine argument).
         for kern in node.kernels():
             KernelModuleInlineTrans().validate(kern)
             ksched = kern.get_kernel_schedule()
             global_variables = ksched.symbol_table.imported_symbols
             if global_variables:
                 raise TransformationError(
                     f"The Symbol Table for kernel '{kern.name}' contains the "
                     f"following symbols with 'global' scope: "
                     f"{[sym.name for sym in global_variables]}. An OpenCL "
                     f"kernel cannot call other kernels and all of the data it "
                     f"accesses must be passed by argument. Use the "
                     f"KernelImportsToArguments transformation to convert such "
                     f"symbols to kernel arguments first.")
  
         # In OpenCL all kernel loops should iterate the whole grid
         for kernel in node.kernels():
             inner_loop = kernel.ancestor(GOLoop)
             outer_loop = inner_loop.ancestor(GOLoop)
             if not (inner_loop.field_space == "go_every" and
                     outer_loop.field_space == "go_every" and
                     inner_loop.iteration_space == "go_all_pts" and
                     outer_loop.iteration_space == "go_all_pts"):
                 raise TransformationError(
                     f"The kernel '{kernel.name}' does not iterate over all "
                     f"grid points. This is a necessary requirement for "
                     f"generating the OpenCL code and can be done by applying "
                     f"the GOMoveIterationBoundariesInsideKernelTrans to each "
                     f"kernel before the GOOpenCLTrans.")
  

References psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans._enable_profiling, psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans._out_of_order, and psyclone.domain.gocean.transformations.gocean_opencl_trans.GOOpenCLTrans._transformed_invokes.

Here is the caller graph for this function:

The documentation for this class was generated from the following file:

/home/docs/checkouts/readthedocs.org/user_builds/psyclone-ref/checkouts/latest/src/psyclone/domain/gocean/transformations/gocean_opencl_trans.py

Public Member Functions

Detailed Description

Member Function Documentation

◆ apply()

◆ name()

◆ validate()