r"""Get CUDA kernel from compiled module.

        Parameters
        ----------
        name : str
            String name of the kernel.
        signature : str
            Function signature for the kernel. For example, if a kernel is
            declared as::

                extern "C" __global__ void axpy(const float *x, double *y, int alpha)

            Then its signature should be::

                const float *x, double *y, int alpha

            or::

                const float *, double *, int

            Note that `*` in signature marks an argument as array and
            `const` marks an argument as constant (input) array.

        Returns
        -------
        CudaKernel
            CUDA kernels that can be launched on GPUs.

def get_kernel(self, name, signature):
        r"""Get CUDA kernel from compiled module.

        Parameters
        ----------
        name : str
            String name of the kernel.
        signature : str
            Function signature for the kernel. For example, if a kernel is
            declared as::

                extern "C" __global__ void axpy(const float *x, double *y, int alpha)

            Then its signature should be::

                const float *x, double *y, int alpha

            or::

                const float *, double *, int

            Note that `*` in signature marks an argument as array and
            `const` marks an argument as constant (input) array.

        Returns
        -------
        CudaKernel
            CUDA kernels that can be launched on GPUs.
        """
        hdl = CudaKernelHandle()
        is_ndarray = []
        is_const = []
        dtypes = []
        pattern = re.compile(r"""^\s*(const)?\s*([\w_]+)\s*(\*)?\s*([\w_]+)?\s*$""")
        args = re.sub(r"\s+", " ", signature).split(",")
        for arg in args:
            match = pattern.match(arg)
            if not match or match.groups()[1] == 'const':
                raise ValueError(
                    'Invalid function prototype "%s". Must be in the '
                    'form of "(const) type (*) (name)"'%arg)
            is_const.append(bool(match.groups()[0]))
            dtype = match.groups()[1]
            is_ndarray.append(bool(match.groups()[2]))
            if dtype not in _DTYPE_CPP_TO_NP:
                raise TypeError(
                    "Unsupported kernel argument type %s. Supported types are: %s."%(
                        arg, ','.join(_DTYPE_CPP_TO_NP.keys())))
            dtypes.append(_DTYPE_NP_TO_MX[_DTYPE_CPP_TO_NP[dtype]])

        check_call(_LIB.MXRtcCudaKernelCreate(
            self.handle,
            c_str(name),
            len(dtypes),
            c_array_buf(ctypes.c_int, array('i', is_ndarray)),
            c_array_buf(ctypes.c_int, array('i', is_const)),
            c_array_buf(ctypes.c_int, array('i', dtypes)),
            ctypes.byref(hdl)))

        return CudaKernel(hdl, name, is_ndarray, dtypes)