How to use the psyclone.f2pygen.CommentGen function in PSyclone

def zero_reduction_variables(red_call_list, parent):
    '''zero all reduction variables associated with the calls in the call
    list'''
    if red_call_list:
        from psyclone.f2pygen import CommentGen
        parent.add(CommentGen(parent, ""))
        parent.add(CommentGen(parent, " Zero summation variables"))
        parent.add(CommentGen(parent, ""))
        for call in red_call_list:
            call.zero_reduction_variable(parent)
        parent.add(CommentGen(parent, ""))

# pylint: disable=too-many-locals
        '''
        Generate code to create data buffers on OpenCL device.

        :param parent: Parent subroutine in f2pygen AST of generated code.
        :type parent: :py:class:`psyclone.f2pygen.SubroutineGen`
        '''
        from psyclone.f2pygen import UseGen, CommentGen, IfThenGen, DeclGen, \
            AssignGen
        grid_arg = self._arguments.find_grid_access()
        symtab = self.root.symbol_table
        # Ensure the fields required by this kernel are on device. We must
        # create the buffers for them if they're not.
        parent.add(UseGen(parent, name="fortcl", only=True,
                          funcnames=["create_rw_buffer"]))
        parent.add(CommentGen(parent, " Ensure field data is on device"))
        for arg in self._arguments.args:
            if arg.type == "field" or \
               (arg.type == "grid_property" and not arg.is_scalar()):
                api_config = Config.get().api_conf("gocean1.0")
                if arg.type == "field":
                    # fields have a 'data_on_device' property for keeping
                    # track of whether they are on the device
                    condition = ".NOT. {0}%data_on_device".format(arg.name)
                    device_buff = "{0}%device_ptr".format(arg.name)
                    host_buff = api_config.grid_properties["go_grid_data"]\
                        .fortran.format(arg.name)
                else:
                    # grid properties do not have such an attribute (because
                    # they are just pointers) so we check whether the device
                    # pointer is NULL.
                    device_buff = "{0}%grid%{1}_device".format(grid_arg.name,

parent.add(DeclGen(parent, datatype="integer", kind="c_size_t",
                                   entity_decls=[nbytes]))
                parent.add(DeclGen(parent, datatype="integer",
                                   kind="c_intptr_t", target=True,
                                   entity_decls=[wevent]))
                api_config = Config.get().api_conf("gocean1.0")
                props = api_config.grid_properties
                num_x = props["go_grid_nx"].fortran.format(grid_arg.name)
                num_y = props["go_grid_ny"].fortran.format(grid_arg.name)
                # Use c_sizeof() on first element of array to be copied over in
                # order to cope with the fact that some grid properties are
                # integer.
                size_expr = "int({0}*{1}, 8)*c_sizeof({2}(1,1))" \
                            .format(num_x, num_y, host_buff)
                ifthen.add(AssignGen(ifthen, lhs=nbytes, rhs=size_expr))
                ifthen.add(CommentGen(ifthen, " Create buffer on device"))
                # Get the name of the list of command queues (set in
                # psyGen.InvokeSchedule)
                qlist = symtab.lookup_with_tag("opencl_cmd_queues").name
                flag = symtab.lookup_with_tag("opencl_error").name

                ifthen.add(AssignGen(ifthen, lhs=device_buff,
                                     rhs="create_rw_buffer(" + nbytes + ")"))
                ifthen.add(
                    AssignGen(ifthen, lhs=flag,
                              rhs="clEnqueueWriteBuffer({0}(1), {1}, CL_TRUE, "
                              "0_8, {2}, C_LOC({3}), 0, C_NULL_PTR, "
                              "C_LOC({4}))".format(qlist, device_buff,
                                                   nbytes, host_buff, wevent)))
                if arg.type == "field":
                    ifthen.add(AssignGen(
                        ifthen, lhs="{0}%data_on_device".format(arg.name),

"get_kernel_by_name"]))
            # Command queues
            parent.add(DeclGen(parent, datatype="integer", save=True,
                               entity_decls=["num_cmd_queues"]))
            parent.add(DeclGen(parent, datatype="integer", save=True,
                               pointer=True, kind="c_intptr_t",
                               entity_decls=["cmd_queues(:)"]))
            parent.add(DeclGen(parent, datatype="integer",
                               entity_decls=["ierr"]))
            parent.add(DeclGen(parent, datatype="logical", save=True,
                               entity_decls=["first_time"],
                               initial_values=[".true."]))
            if_first = IfThenGen(parent, "first_time")
            parent.add(if_first)
            if_first.add(AssignGen(if_first, lhs="first_time", rhs=".false."))
            if_first.add(CommentGen(if_first,
                                    " Ensure OpenCL run-time is initialised "
                                    "for this PSy-layer module"))
            if_first.add(CallGen(if_first, "psy_init"))
            if_first.add(AssignGen(if_first, lhs="num_cmd_queues",
                                   rhs="get_num_cmd_queues()"))
            if_first.add(AssignGen(if_first, lhs="cmd_queues", pointer=True,
                                   rhs="get_cmd_queues()"))
            # Kernel pointers
            kernels = self.walk(self._children, Call)
            for kern in kernels:
                kernel = "kernel_" + kern.name  # TODO use namespace manager
                parent.add(
                    DeclGen(parent, datatype="integer", kind="c_intptr_t",
                            save=True, target=True, entity_decls=[kernel]))
                if_first.add(
                    AssignGen(

# Add a callback here so that derived classes can adjust the list
        # of variables to provide, or the suffix used (which might
        # depend on the variable name which could create clashes).
        self.update_vars_and_postname()

        options = {'pre_var_list': self._input_list,
                   'post_var_list': self._output_list,
                   'post_var_postfix': self._post_name}

        from psyclone.f2pygen import CommentGen
        parent.add(CommentGen(parent, ""))
        parent.add(CommentGen(parent, " ExtractStart"))
        parent.add(CommentGen(parent, ""))
        super(ExtractNode, self).gen_code(parent, options)
        parent.add(CommentGen(parent, ""))
        parent.add(CommentGen(parent, " ExtractEnd"))
        parent.add(CommentGen(parent, ""))

# value for "go_grid_data" is actually used.
        from psyclone.psyGen import CodedKern
        for kernel in self.psy_data_body.walk(CodedKern):
            kernel.clear_cached_data()

        # Recreate the instrumented region. Due to the changes in the
        # config files, fields and properties will now become local
        # plain arrays and variables:
        for child in self.psy_data_body:
            child.gen_code(prog)

        # Now reset all properties back to the original values:
        for name in all_props.keys():
            all_props[name] = orig_props[name]

        prog.add(CommentGen(prog, " RegionEnd"))
        prog.add(CommentGen(prog, ""))

        for var_name in output_list:
            prog.add(CommentGen(prog, " Check {0}".format(var_name)))

        code = str(module.root)

        with open("driver-{0}-{1}.f90".
                  format(module_name, region_name), "w") as out:
            out.write(code)

sub.add(UseGen(sub, name="fortcl", only=True,
                       funcnames=["ocl_env_init", "add_kernels"]))
        # Add a logical variable used to ensure that this routine is only
        # executed once.
        sub.add(DeclGen(sub, datatype="logical", save=True,
                        entity_decls=["initialised"],
                        initial_values=[".False."]))
        # Check whether or not this is our first time in the routine
        sub.add(CommentGen(sub, " Check to make sure we only execute this "
                           "routine once"))
        ifthen = IfThenGen(sub, ".not. initialised")
        sub.add(ifthen)
        ifthen.add(AssignGen(ifthen, lhs="initialised", rhs=".True."))

        # Initialise the OpenCL environment
        ifthen.add(CommentGen(ifthen,
                              " Initialise the OpenCL environment/device"))
        ifthen.add(CallGen(ifthen, "ocl_env_init"))

        # Create a list of our kernels
        ifthen.add(CommentGen(ifthen,
                              " The kernels this PSy layer module requires"))
        nkernstr = str(len(kernels))

        # Declare array of character strings
        ifthen.add(CharDeclGen(
            ifthen, length="30",
            entity_decls=["kernel_names({0})".format(nkernstr)]))
        for idx, kern in enumerate(kernels):
            ifthen.add(AssignGen(ifthen, lhs="kernel_names({0})".format(idx+1),
                                 rhs='"{0}"'.format(kern)))
        ifthen.add(CommentGen(ifthen,

qlist = symtab.lookup_with_tag("opencl_cmd_queues").name
        flag = symtab.lookup_with_tag("opencl_error").name

        # Then we call clEnqueueNDRangeKernel
        parent.add(CommentGen(parent, " Launch the kernel"))
        cnull = "C_NULL_PTR"
        queue_number = self._opencl_options['queue_number']
        cmd_queue = qlist + "({0})".format(queue_number)

        args = ", ".join([cmd_queue, kernel, "2", cnull,
                          "C_LOC({0})".format(glob_size),
                          "C_LOC({0})".format(local_size),
                          "0", cnull, cnull])
        parent.add(AssignGen(parent, lhs=flag,
                             rhs="clEnqueueNDRangeKernel({0})".format(args)))
        parent.add(CommentGen(parent, ""))