class_docstr: Compile and run CUDA code from Python.
  
  In CUDA 7.5, you need to prepend your kernel definitions
  with 'extern "C"' to avoid name mangling::
  
      source = r'''
      extern "C" __global__ void axpy(const float *x, float *y, float alpha) {
          int i = threadIdx.x + blockIdx.x * blockDim.x;
          y[i] += alpha * x[i];
      }
      '''
      module = mx.rtc.CudaModule(source)
      func = module.get_kernel("axpy", "const float *x, float *y, float alpha")
      x = mx.nd.ones((10,), ctx=mx.gpu(0))
      y = mx.nd.zeros((10,), ctx=mx.gpu(0))
      func.launch([x, y, 3.0], mx.gpu(0), (1, 1, 1), (10, 1, 1))
      print(y)
  
  Starting from CUDA 8.0, you can instead export functions by name.
  This also allows you to use templates::
  
      source = r'''
      template<typename DType>
      __global__ void axpy(const DType *x, DType *y, DType alpha) {
          int i = threadIdx.x + blockIdx.x * blockDim.x;
          y[i] += alpha * x[i];
      }
      '''
      module = mx.rtc.CudaModule(source, exports=['axpy<float>', 'axpy<double>'])
      func32 = module.get_kernel("axpy<float>", "const float *x, float *y, float alpha")
      x = mx.nd.ones((10,), dtype='float32', ctx=mx.gpu(0))
      y = mx.nd.zeros((10,), dtype='float32', ctx=mx.gpu(0))
      func32.launch([x, y, 3.0], mx.gpu(0), (1, 1, 1), (10, 1, 1))
      print(y)
  
      func64 = module.get_kernel("axpy<double>", "const double *x, double *y, double alpha")
      x = mx.nd.ones((10,), dtype='float64', ctx=mx.gpu(0))
      y = mx.nd.zeros((10,), dtype='float64', ctx=mx.gpu(0))
      func32.launch([x, y, 3.0], mx.gpu(0), (1, 1, 1), (10, 1, 1))
      print(y)
  
  
  Parameters
  ----------
  source : str
      Complete source code.
  options : tuple of str
      Compiler flags. For example, use "-I/usr/local/cuda/include" to
      add cuda headers to include path.
  exports : tuple of str
      Export kernel names.
class_docstr_tok: ["Compile","and","run","CUDA","code","from","Python",".","In","CUDA","7.5",",","you","need","to","prepend","your","kernel","definitions","with","'extern","``","C","''","'","to","avoid","name","mangling",":",":","source","=","r","''","'","extern","``","C","''","__global__","void","axpy","(","const","float","*","x",",","float","*","y",",","float","alpha",")","{","int","i","=","threadIdx.x","+","blockIdx.x","*","blockDim.x",";","y","[","i","]","+=","alpha","*","x","[","i","]",";","}","``","'","module","=","mx.rtc.CudaModule","(","source",")","func","=","module.get_kernel","(","``","axpy","''",",","``","const","float","*","x",",","float","*","y",",","float","alpha","''",")","x","=","mx.nd.ones","(","(","10",",",")",",","ctx=mx.gpu","(","0",")",")","y","=","mx.nd.zeros","(","(","10",",",")",",","ctx=mx.gpu","(","0",")",")","func.launch","(","[","x",",","y",",","3.0","]",",","mx.gpu","(","0",")",",","(","1",",","1",",","1",")",",","(","10",",","1",",","1",")",")","print","(","y",")","Starting","from","CUDA","8.0",",","you","can","instead","export","functions","by","name",".","This","also","allows","you","to","use","templates",":",":","source","=","r","''","'","template","<","typename","DType",">","__global__","void","axpy","(","const","DType","*","x",",","DType","*","y",",","DType","alpha",")","{","int","i","=","threadIdx.x","+","blockIdx.x","*","blockDim.x",";","y","[","i","]","+=","alpha","*","x","[","i","]",";","}","``","'","module","=","mx.rtc.CudaModule","(","source",",","exports=","[","'axpy","<","float",">","'",",","'axpy","<","double",">","'","]",")","func32","=","module.get_kernel","(","``","axpy","<","float",">","''",",","``","const","float","*","x",",","float","*","y",",","float","alpha","''",")","x","=","mx.nd.ones","(","(","10",",",")",",","dtype='float32","'",",","ctx=mx.gpu","(","0",")",")","y","=","mx.nd.zeros","(","(","10",",",")",",","dtype='float32","'",",","ctx=mx.gpu","(","0",")",")","func32.launch","(","[","x",",","y",",","3.0","]",",","mx.gpu","(","0",")",",","(","1",",","1",",","1",")",",","(","10",",","1",",","1",")",")","print","(","y",")","func64","=","module.get_kernel","(","``","axpy","<","double",">","''",",","``","const","double","*","x",",","double","*","y",",","double","alpha","''",")","x","=","mx.nd.ones","(","(","10",",",")",",","dtype='float64","'",",","ctx=mx.gpu","(","0",")",")","y","=","mx.nd.zeros","(","(","10",",",")",",","dtype='float64","'",",","ctx=mx.gpu","(","0",")",")","func32.launch","(","[","x",",","y",",","3.0","]",",","mx.gpu","(","0",")",",","(","1",",","1",",","1",")",",","(","10",",","1",",","1",")",")","print","(","y",")","Parameters","--","--","--","--","--","source",":","str","Complete","source","code",".","options",":","tuple","of","str","Compiler","flags",".","For","example",",","use","``","-I\/usr\/local\/cuda\/include","''","to","add","cuda","headers","to","include","path",".","exports",":","tuple","of","str","Export","kernel","names","."]
docstr: 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.
docstr_tok: ["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","."]
code: 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"""^(const)?\s?([\w_]+)\s?(\*)?\s?([\w_]+)?$"""
      )
      args = re.sub(r"\s+", " ", signature).split(",")
      for arg in args:
          sanitized_arg = " ".join(arg.split())
          match = pattern.match(sanitized_arg)
          if not match or match.groups()[1] == "const":
              raise ValueError(
                  f'Invalid function prototype "{sanitized_arg}". Must be in the '
                  'form of "(const) type (*) (name)"'
              )
          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 {}. Supported types are: {}.".format(
                      sanitized_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)
  
file_path: mxnet/python/mxnet/rtc.py
imports: from array import array
  import re
  import ctypes
  import numpy
  from .base import _LIB, mx_uint, c_array, c_array_buf, c_str_array, check_call
  from .base import c_str, CudaModuleHandle, CudaKernelHandle, numeric_types, string_types
  from .ndarray import dtype_np_to_mx, dtype_mx_to_np, NDArray
label_desc: Use image node_id 3 for calling the CudaModule obj's underlying member method code with example usage: obj.get_kernel(name, signature) and returns: CudaKernel