Inheritance diagram for psyclone.psyir.backend.opencl.OpenCLWriter:

[legend]

Collaboration diagram for psyclone.psyir.backend.opencl.OpenCLWriter:

[legend]

Public Member Functions
def	__init__ (self, skip_nodes=False, indent_string=" ", initial_indent_depth=0, kernels_local_size=1)

def	gen_id_variable (self, symbol, dimension_index)

def	gen_declaration (self, symbol)

def	gen_array_length_variables (self, symbol, symtab=None)

def	kernelschedule_node (self, node)

Public Member Functions inherited from psyclone.psyir.backend.c.CWriter
def	gen_indices (self, indices, var_name=None)

def	gen_local_variable (self, symbol)

def	assignment_node (self, node)

def	literal_node (self, node)

def	ifblock_node (self, node)

def	unaryoperation_node (self, node)

def	binaryoperation_node (self, node)

def	intrinsiccall_node (self, node)

def	return_node (self, _)

def	codeblock_node (self, _)

def	loop_node (self, node)

def	regiondirective_node (self, node)

def	standalonedirective_node (self, node)

def	filecontainer_node (self, node)

Public Member Functions inherited from psyclone.psyir.backend.language_writer.LanguageWriter
def	__init__ (self, array_parenthesis, structure_character, skip_nodes=False, indent_string=" ", initial_indent_depth=0, check_global_constraints=True)

def	array_parenthesis (self)

def	structure_character (self)

def	arrayreference_node (self, node)

def	structurereference_node (self, node)

def	member_node (self, node)

def	arrayofstructuresreference_node (self, node)

def	clause_node (self, node)

Public Member Functions inherited from psyclone.psyir.backend.visitor.PSyIRVisitor
def	reference_node (self, node)

def	__call__ (self, node)

Detailed Description

Implements a PSyIR-to-OpenCL back-end for the PSyIR AST. This writer
produces OpenCL code conforming to Version 1.2 of the specification
(https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf).

:param bool skip_nodes: if skip_nodes is False then an exception \
is raised if a visitor method for a PSyIR node has not been \
implemented, otherwise the visitor silently continues. This is an \
optional argument which defaults to False.
:param str indent_string: specifies what to use for indentation. This \
is an optional argument that defaults to two spaces.
:param int initial_indent_depth: specifies how much indentation to \
start with. This is an optional argument that defaults to 0.
:param int kernel_local_size: uses the given local_size when generating \
OpenCL kernels.

:raises TypeError: if kernel_local_size is not an integer.
:raises ValueError: if kernel_local_size is not positive.

Definition at line 48 of file opencl.py.

Member Function Documentation

◆ gen_array_length_variables()

def psyclone.psyir.backend.opencl.OpenCLWriter.gen_array_length_variables	(	self,
		symbol,
		symtab = `None`
	)

Generate length variable declaration and initialisation for each array
dimension of the given symbol. By convention they are named:
<name>LEN<DIM>, and initialised using the function:
`size_t get_global_size(uint dimindx)`.


:param symbol: The symbol instance.
:type symbol: :py:class:`psyclone.psyir.symbols.DataSymbol`
:param symtab: The symbol table from the given symbol context to \
    check for name clashes.
:type symtab: :py:class:`psyclone.psyir.symbols.SymbolTable`

:return: OpenCL declaration and initialisation of length variables.
:rtype: str

:raises VisitorError: if the array length variable name clashes \
    with another symbol name.

Definition at line 150 of file opencl.py.

     def gen_array_length_variables(self, symbol, symtab=None):
         '''
         Generate length variable declaration and initialisation for each array
         dimension of the given symbol. By convention they are named:
         <name>LEN<DIM>, and initialised using the function:
         `size_t get_global_size(uint dimindx)`.
  
  
         :param symbol: The symbol instance.
         :type symbol: :py:class:`psyclone.psyir.symbols.DataSymbol`
         :param symtab: The symbol table from the given symbol context to \
             check for name clashes.
         :type symtab: :py:class:`psyclone.psyir.symbols.SymbolTable`
  
         :return: OpenCL declaration and initialisation of length variables.
         :rtype: str
  
         :raises VisitorError: if the array length variable name clashes \
             with another symbol name.
         '''
  
         code = ""
         dimensions = len(symbol.shape)
         for dim in range(1, dimensions + 1):
             code += self._nindent + "int "
             varname = symbol.name + "LEN" + str(dim)
  
             # Check there is no clash with other variables
             if symtab and varname in symtab:
                 raise VisitorError(
                     f"Unable to declare the variable '{varname}' to store the "
                     f"length of '{symbol.name}' because the Symbol Table "
                     f"already contains a symbol with the same name.")
  
             code += varname + " = get_global_size("
             code += str(dim - 1) + ");\n"
         return code
  

References psyclone.psyir.backend.visitor.PSyIRVisitor._nindent().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ gen_declaration()

def psyclone.psyir.backend.opencl.OpenCLWriter.gen_declaration	(	self,
		symbol
	)

Generate the declaration of an OpenCL variable. This can be either
a local variable or a routine argument, so no indention or punctuation
is generated by this method.

Memory buffers reside in specific levels of the memory hierarchy,
and pointers are annotated with the region qualifiers __global,
__local, __constant, and __private, reflecting this.
At the moment all memory buffers (PSyIR arrays) are allocated at
the global address space.

:param symbol: The symbol instance.
:type symbol: :py:class:`psyclone.psyir.symbols.DataSymbol`

:returns: The OpenCL declaration of the given of the symbol.
:rtype: str

:raises VisitorError: if an array is encountered that does not have \
                      a lower bound of 1 for all of its dimensions.

Reimplemented from psyclone.psyir.backend.c.CWriter.

Definition at line 113 of file opencl.py.

     def gen_declaration(self, symbol):
         '''
         Generate the declaration of an OpenCL variable. This can be either
         a local variable or a routine argument, so no indention or punctuation
         is generated by this method.
  
         Memory buffers reside in specific levels of the memory hierarchy,
         and pointers are annotated with the region qualifiers __global,
         __local, __constant, and __private, reflecting this.
         At the moment all memory buffers (PSyIR arrays) are allocated at
         the global address space.
  
         :param symbol: The symbol instance.
         :type symbol: :py:class:`psyclone.psyir.symbols.DataSymbol`
  
         :returns: The OpenCL declaration of the given of the symbol.
         :rtype: str
  
         :raises VisitorError: if an array is encountered that does not have \
                               a lower bound of 1 for all of its dimensions.
         '''
         prefix = ""
         if symbol.shape:
             for dim in symbol.shape:
                 if not isinstance(dim, ArrayType.ArrayBounds):
                     continue
                 if (not isinstance(dim.lower, Literal) or
                         dim.lower.value != "1"):
                     raise VisitorError(
                         f"The OpenCL backend only supports arrays with a lower"
                         f" bound of 1 in each dimension. However, array "
                         f"'{symbol.name}' has a lower bound of "
                         f"'{self._visit(dim.lower)}' for dimension "
                         f"{symbol.shape.index(dim)}")
             prefix += "__global "
         return prefix + super(OpenCLWriter, self).gen_declaration(symbol)
  

Here is the caller graph for this function:

◆ gen_id_variable()

def psyclone.psyir.backend.opencl.OpenCLWriter.gen_id_variable	(	self,
		symbol,
		dimension_index
	)

Generate the declaration and initialisation of a variable identifying
an OpenCL work-item. IDs are initialised by the OpenCL function:
`size_t get_global_id(uint dimindx)`

:param symbol: The symbol instance.
:type symbol: :py:class:`psyclone.psyir.symbols.DataSymbol`
:param int dimension_index: Dimension which the given symbol will \
    iterate on.

:returns: OpenCL declaration of an OpenCL id variable.
:rtype: str

:raises VisitorError: if symbol is not a scalar integer

Definition at line 86 of file opencl.py.

     def gen_id_variable(self, symbol, dimension_index):
         '''
         Generate the declaration and initialisation of a variable identifying
         an OpenCL work-item. IDs are initialised by the OpenCL function:
         `size_t get_global_id(uint dimindx)`
  
         :param symbol: The symbol instance.
         :type symbol: :py:class:`psyclone.psyir.symbols.DataSymbol`
         :param int dimension_index: Dimension which the given symbol will \
             iterate on.
  
         :returns: OpenCL declaration of an OpenCL id variable.
         :rtype: str
  
         :raises VisitorError: if symbol is not a scalar integer
         '''
         if (not isinstance(symbol.datatype, ScalarType) or
                 symbol.datatype.intrinsic != ScalarType.Intrinsic.INTEGER):
             raise VisitorError(
                 f"OpenCL work-item identifiers must be scalar integer symbols "
                 f"but found {symbol}.")
  
         code = ""
         code += self._nindent + "int " + symbol.name
         code += " = get_global_id(" + str(dimension_index) + ");\n"
         return code
  

References psyclone.psyir.backend.visitor.PSyIRVisitor._nindent().

Here is the call graph for this function:

Here is the caller graph for this function:

◆ kernelschedule_node()

def psyclone.psyir.backend.opencl.OpenCLWriter.kernelschedule_node	(	self,
		node
	)

This method is called when a KernelSchedule instance is found in
the PSyIR tree.

:param node: A KernelSchedule PSyIR node.
:type node: :py:class:`psyclone.psyir.nodes.KernelSchedule`

:returns: The OpenCL code as a string.
:rtype: str

:raises VisitorError: if a non-precision symbol is found with a \
                      UnresolvedInterface.

Definition at line 188 of file opencl.py.

     def kernelschedule_node(self, node):
         '''This method is called when a KernelSchedule instance is found in
         the PSyIR tree.
  
         :param node: A KernelSchedule PSyIR node.
         :type node: :py:class:`psyclone.psyir.nodes.KernelSchedule`
  
         :returns: The OpenCL code as a string.
         :rtype: str
  
         :raises VisitorError: if a non-precision symbol is found with a \
                               UnresolvedInterface.
         '''
         # OpenCL implementation assumptions:
         # - All array have the same size and it is given by the
         #   global_work_size argument to clEnqueueNDRangeKernel.
         # - Assumes no dependencies among kernels called concurrently.
         # - All real variables are 64-bit
  
         # TODO: At the moment, the method caller is responsible to ensure
         # these assumptions. KernelSchedule access to the kernel
         # meta-arguments could be used to check them and also improve the
         # generated code. (Issue #288)
  
         symtab = node.symbol_table
         data_args = symtab.data_arguments
  
         # Check that we know where everything in the symbol table
         # comes from.  TODO #592 ultimately precision symbols should
         # be included in this check too as we will need to be able to
         # map from them to the equivalent OpenCL type.
         unresolved_datasymbols = list(set(symtab.unresolved_datasymbols) -
                                       set(symtab.precision_datasymbols))
         if unresolved_datasymbols:
             symbols_txt = ", ".join(
                 [f"'{sym.name}'" for sym in unresolved_datasymbols])
             raise VisitorError(
                 f"Cannot generate OpenCL because the symbol table contains "
                 f"unresolved data entries (i.e. that have no defined "
                 f"Interface) which are not used purely to define the "
                 f"precision of other symbols: {symbols_txt}")
  
         # Start OpenCL kernel definition
         code = self._nindent
         if self._kernels_local_size != 1:
             code += f"__attribute__((reqd_work_group_size("\
                     f"{self._kernels_local_size}, 1, 1)))\n"
         code += "__kernel void " + node.name + "(\n"
         self._depth += 1
         arguments = []
  
         # Declare kernel arguments
         for symbol in data_args:
             arguments.append(self._nindent + self.gen_declaration(symbol))
         code += ",\n".join(arguments) + "\n"
         code += self._nindent + "){\n"
  
         # Declare local variables.
         for symbol in symtab.automatic_datasymbols:
             code += self.gen_local_variable(symbol)
  
         # Declare array length
         for symbol in data_args:
             code += self.gen_array_length_variables(symbol, symtab)
  
         # Declare iteration indices
         for index, symbol in enumerate(symtab.iteration_indices):
             code += self.gen_id_variable(symbol, index)
  
         # Generate kernel body
         for child in node.children:
             code += self._visit(child)
  
         # Close kernel definition
         self._depth -= 1
         code += self._nindent + "}\n\n"
  
         return code

References psyclone.psyir.backend.visitor.PSyIRVisitor._depth, psyclone.psyir.backend.opencl.OpenCLWriter._kernels_local_size, psyclone.psyir.backend.visitor.PSyIRVisitor._nindent(), psyclone.psyir.backend.visitor.PSyIRVisitor._visit(), psyclone.psyir.backend.opencl.OpenCLWriter.gen_array_length_variables(), psyclone.psyir.backend.c.CWriter.gen_declaration(), psyclone.psyir.backend.opencl.OpenCLWriter.gen_declaration(), psyclone.psyir.backend.opencl.OpenCLWriter.gen_id_variable(), and psyclone.psyir.backend.c.CWriter.gen_local_variable().

Here is the call 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/psyir/backend/opencl.py

Public Member Functions

Detailed Description

Member Function Documentation

◆ gen_array_length_variables()

◆ gen_declaration()

◆ gen_id_variable()

◆ kernelschedule_node()