Add custom opencl

Add a custom operator implementation on OpenCL Backend.

• User can give the OpenCL code snippet to be used within an OpenCL ElementwiseKernel.

hysop/backend/device/opencl/operator/custom.py

+from hysop.constants import DirectionLabels
+from hysop.tools.decorators import debug
+from hysop.core.memory.memory_request import MemoryRequest
+from hysop.backend.device.opencl.opencl_operator import OpenClOperator, op_apply
+from hysop.backend.device.opencl.autotunable_kernels.custom_symbolic import OpenClAutotunableCustomSymbolicKernel
+from hysop.backend.device.opencl.opencl_kernel_launcher import OpenClKernelListLauncher
+from hysop.backend.device.opencl.opencl_copy_kernel_launchers import OpenClCopyBufferRectLauncher
+from hysop.operator.base.custom import CustomOperatorBase
+from hysop.parameters.scalar_parameter import ScalarParameter
+
+from hysop.backend.device.codegen.structs.mesh_info import MeshBaseStruct, MeshInfoStruct
+from hysop.backend.device.codegen.base.variables import dtype_to_ctype
+from hysop.backend.device.codegen.base.opencl_codegen import OpenClCodeGenerator
+
+from pyopencl.elementwise import ElementwiseKernel
+import pyopencl as cl
+
+
+class OpenClCustomOperator(CustomOperatorBase, OpenClOperator):
+
+    @debug
+    def __new__(cls, **kwds):
+        return super().__new__(cls, **kwds)
+
+    @debug
+    def __init__(self, **kwds):
+        super().__init__(**kwds)
+
+    @debug
+    def setup(self, *args, **kwargs):
+        super().setup(*args, **kwargs)
+        dim = self.domain.dim
+        cg = OpenClCodeGenerator('test_generator', self.typegen)
+        mbs = MeshBaseStruct(self.typegen, typedef='{}MeshBase{}D_s'.format(self.typegen.fbtype[0], dim),
+                             vsize=dim)
+        cg.require('mis', MeshInfoStruct(self.typegen,
+                                         typedef='{}MeshInfo{}D_s'.format(self.typegen.fbtype[0], dim),
+                                         mbs_typedef=mbs.typedef, vsize=dim))
+        for f in self.discrete_fields:
+            fn = f.continuous_fields()[0].name
+            mesh_info = MeshInfoStruct.create_from_mesh(fn+"_mesh", self.typegen, f.mesh,
+                                                        storage=OpenClCodeGenerator.default_keywords['constant'])
+            mesh_info[1].declare(cg, _const=True)
+            cg.append(f"""int3 get_{fn}i_xyz(int i) {{
+              int iz = i/({fn}_mesh.local_mesh.resolution.x*{fn}_mesh.local_mesh.resolution.y);
+              int iy = (i-({fn}_mesh.local_mesh.resolution.x*{fn}_mesh.local_mesh.resolution.y)*iz)/({fn}_mesh.local_mesh.resolution.x);
+              return (int3)(i % {fn}_mesh.local_mesh.resolution.x,iy,iz);}}""")
+        kernel_args = []
+        for f in self.dinvar:
+            if f not in self.doutvar:
+                fn = f.continuous_fields()[0].name
+                kernel_args.append(f"const {self.typegen.fbtype} * {fn}")
+        for p in self.dinparam:
+            if isinstance(p, ScalarParameter):
+                kernel_args.append(f"const {self.typegen.fbtype} {p.name}")
+                self._cl_dinparam.append(p)
+            else:
+                kernel_args.append(f"const {self.typegen.fbtype} *{p.name}")
+        for f in self.doutvar:
+            fn = f.continuous_fields()[0].name
+            kernel_args.append(f"{self.typegen.fbtype} * {fn}")
+        for p in self.doutparam:
+            kernel_args.append(f"{self.typegen.fbtype} *{p.name}")
+        self.__elementwise = ElementwiseKernel(
+            self.cl_env.context,
+            ",".join(kernel_args),
+            self.func,
+            f"__{self.name}_elementwise", preamble=str(cg))
+
+        # Build testing:
+        try:
+            self.__elementwise.get_kernel(False)
+        except cl.RuntimeError as e:
+            print("USED KERNEL")
+            print(",".join(kernel_args))
+            print(str(cg)+self.func)
+            raise e
+
+    @op_apply
+    def apply(self, **kwds):
+        super().apply(**kwds)
+        args = (tuple(_.sbuffer for _ in self.dinvar)
+                + tuple(cl.array.to_device(self.cl_env.default_queue, _._value) for _ in self.dinparam)
+                + tuple(_.sbuffer for _ in self.doutvar)
+                + tuple(cl.array.to_device(self.cl_env.default_queue, _._value) for _ in self.doutparam)
+                + tuple(cl.array.to_device(self.cl_env.default_queue, _._value) for _ in self.extra_args))
+        self.__elementwise(*args)
+        for gh_exch in self.ghost_exchanger:
+            gh_exch.exchange_ghosts()