class_docstr: Parameters
  ----------
  axis
      Can not be None. If you want the array flattened, do it before.
class_docstr_tok: ["Parameters","--","--","--","--","--","axis","Can","not","be","None",".","If","you","want","the","array","flattened",",","do","it","before","."]
code: def gpu_kernels(self, node, nodename):
      kernels = []
      # cumadd
      kname = "k_cumadd"
      op = {"mul": "*", "add": "+"}[self.mode]
      k_var = "k_cumadd_" + nodename
      dtype_x = node.inputs[0].dtype
      flags = Kernel.get_flags(dtype_x)
      code = (
          """#include "cluda.h"
  
          KERNEL void %(kname)s(float* input, ga_size input_offset,
                                float* output, ga_size output_offset,
                                ga_ssize inputStrides_x, ga_ssize inputStrides_y, ga_ssize inputStrides_z,
                                ga_ssize outputStrides_x, ga_ssize outputStrides_y, ga_ssize outputStrides_z,
                                const int offsetY, const int offsetZ,
                                const int beforeLastElementIdx, const int lastElementIdx){
              input = (float *)(((char *)input) + input_offset);
              output = (float *)(((char *)output) + output_offset);
              int idY = blockIdx.y + offsetY;
              int idZ = blockIdx.z + offsetZ;
  
              int dataOffsetY_input = idY * inputStrides_y + idZ * inputStrides_z;
              int dataOffsetY_output = idY * outputStrides_y + idZ * outputStrides_z;
              int idx_last_input = lastElementIdx*inputStrides_x + dataOffsetY_input;
              int idx_last_output = lastElementIdx*outputStrides_x + dataOffsetY_output;
              int idx_beforelast = beforeLastElementIdx*outputStrides_x + dataOffsetY_output;
              output[idx_last_output] = input[idx_last_input] %(op)s output[idx_beforelast];
              }
          """
          % locals()
      )
      params = [
          gpuarray.GpuArray,
          gpuarray.SIZE,
          gpuarray.GpuArray,
          gpuarray.SIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          "intc",
          "intc",
          "intc",
          "intc",
      ]
      kernels.append(
          Kernel(
              code=code,
              name=kname,
              params=params,
              flags=flags,
              objvar=k_var,
          )
      )
      # blockCumOp
      kname = "k_blockCumOp"
      k_var = "k_blockCumOp_" + nodename
      params = [
          gpuarray.GpuArray,
          gpuarray.SIZE,
          gpuarray.GpuArray,
          gpuarray.SIZE,
          gpuarray.SIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          "int32",
          "int32",
          gpuarray.GpuArray,
          gpuarray.SIZE,
      ]
      code = (
          """#include "cluda.h"
  
          // helper functions
          WITHIN_KERNEL
          void k_reductionPhase(float* partialCumOp) {
              // Traverse down from leaves to root building partial sums at internal nodes in the tree.
              for (unsigned int stride = 1; stride <= blockDim.x; stride *= 2) {
                  local_barrier();
                  unsigned int index = (threadIdx.x + 1) * (stride * 2) - 1;
                  if (index < blockDim.x*2) {
                      partialCumOp[index] %(op)s= partialCumOp[index - stride];
                  }
              }
          }
  
          WITHIN_KERNEL
          void k_fetchData(float* partialCumOp, float* input, int globalThreadID,
                           ga_ssize dataStrides_x, ga_ssize dataStrides_y, ga_ssize dataStrides_z,
                           int offsetY, int offsetZ) {
              // blockIdx.y and blockIdx.z represents the current independent cum op
              int idY = blockIdx.y + offsetY;
              int idZ = blockIdx.z + offsetZ; int offset = idY * dataStrides_y + idZ * dataStrides_z;
              int idx_even = (globalThreadID*2    ) * dataStrides_x + offset;
              int idx_odd  = (globalThreadID*2 + 1) * dataStrides_x + offset;
              partialCumOp[threadIdx.x*2]     = input[idx_even];
              partialCumOp[threadIdx.x*2 + 1] = input[idx_odd];
          }
  
          WITHIN_KERNEL
          void k_reversePhase(float* partialCumOp) {
              // Traverse back up the tree building the scan from the partial sums
              for (unsigned int stride = exp2(ceil(log2((float)blockDim.x))); stride > 0; stride /= 2) {
                  local_barrier();
                  unsigned int index = (threadIdx.x + 1) * (stride * 2) - 1;
                  if (index + stride < blockDim.x*2) {
                      partialCumOp[index + stride] %(op)s= partialCumOp[index];
                  }
              }
          }
  
          WITHIN_KERNEL
          void k_pushData(float* partialCumOp, float* output, int globalThreadID,
                          ga_ssize dataStrides_x, ga_ssize dataStrides_y, ga_ssize dataStrides_z,
                          int offsetY, int offsetZ) {
              local_barrier();
              // blockIdx.y and blockIdx.z represents the current independent cum op
              int idY = blockIdx.y + offsetY;
              int idZ = blockIdx.z + offsetZ;
              int offset = idY * dataStrides_y + idZ * dataStrides_z;
              int idx_even = (globalThreadID*2    ) * dataStrides_x + offset;
              int idx_odd  = (globalThreadID*2 + 1) * dataStrides_x + offset;
              output[idx_even] = partialCumOp[threadIdx.x*2];
              output[idx_odd]  = partialCumOp[threadIdx.x*2 + 1];
          }
  
          KERNEL void k_blockCumOp(float* input, ga_size input_offset,
                                   float* output, ga_size output_offset,
                                   size_t nbElementsPerCumOp, ga_ssize inputStrides_x,
                                   ga_ssize inputStrides_y,  ga_ssize inputStrides_z,
                                   ga_ssize outputStrides_x, ga_ssize outputStrides_y,
                                   ga_ssize outputStrides_z, int offsetY,
                                   int offsetZ, float* blockSum, ga_size blockSum_offset) {
              input = (float *)(((char *)input) + input_offset);
              output = (float *)(((char *)output) + output_offset);
              blockSum = (float *)(((char *)blockSum) + blockSum_offset);
  
              // Regarding blockIdx and threadIdx, 'CumOp' is always performed along the X axis.
              // The Y and Z axis of the grid will contain all independent cumops of the 2D/3D case.
  
              int globalThreadID = blockIdx.x * blockDim.x + threadIdx.x;
  
              // Check if current thread has data to process.
              if (globalThreadID >= (nbElementsPerCumOp+1)/2) {
                  return;
              }
  
              extern __shared__ float partialCumOp[];
  
              // Load data in shared memory
              k_fetchData(partialCumOp, input, globalThreadID, inputStrides_x, inputStrides_y, inputStrides_z, offsetY, offsetZ);
  
              // Use a dichotomy approach to compute the cum op (i.e. balanced binary tree).
              // The tree is sweeped from the leaves to the root and from the root to the leaves.
              // Similar to http://www.umiacs.umd.edu/~ramani/cmsc828e_gpusci/ScanTalk.pdf
              k_reductionPhase(partialCumOp);
              k_reversePhase(partialCumOp);
  
              // Write the final output to global memory
              k_pushData(partialCumOp, output, globalThreadID, outputStrides_x, outputStrides_y, outputStrides_z, offsetY, offsetZ);
  
              if (blockSum != NULL){
                  if (threadIdx.x == blockDim.x - 1) {
                      blockSum[blockIdx.x*(gridDim.y*gridDim.z) + (blockIdx.y + offsetY)*gridDim.z + blockIdx.z + offsetZ] = partialCumOp[threadIdx.x*2 + 1];
                  }
              }
          }
          """
          % locals()
      )
      kernels.append(
          Kernel(
              code=code,
              name=kname,
              params=params,
              flags=flags,
              objvar=k_var,
          )
      )
      # k_finalCumOp
      kname = "k_finalCumOp"
      k_var = "k_finalCumOp_" + nodename
      code = (
          """#include "cluda.h"
  
          KERNEL void k_finalCumOp(float* output, ga_size output_offset,
                                   float* blockSum, ga_size blockSum_offset,
                                   size_t nbElementsPerCumOp,
                                   ga_ssize dataStrides_x,  ga_ssize dataStrides_y,  ga_ssize dataStrides_z,
                                   int offsetY, int offsetZ) {
  
              output = (float *)(((char *)output) + output_offset);
              blockSum = (float *)(((char *)blockSum) + blockSum_offset);
  
              int globalThreadID = (blockIdx.x + 1) * blockDim.x + threadIdx.x;
  
              // Check if current has data to process.
              if (globalThreadID >= (nbElementsPerCumOp+1)/2)
                  return;
  
              int idY = blockIdx.y + offsetY;
              int idZ = blockIdx.z + offsetZ;
  
              const float currentBlockSum = blockSum[blockIdx.x*(gridDim.y*gridDim.z) + idY*gridDim.z + idZ];
  
              int offset = idY * dataStrides_y + idZ * dataStrides_z;
              int idx_even = (globalThreadID*2    ) * dataStrides_x + offset;
              int idx_odd  = (globalThreadID*2 + 1) * dataStrides_x + offset;
              output[idx_even] %(op)s= currentBlockSum;
              output[idx_odd] %(op)s= currentBlockSum;
          }
          """
          % locals()
      )
      params = [
          gpuarray.GpuArray,
          gpuarray.SIZE,
          gpuarray.GpuArray,
          gpuarray.SIZE,
          gpuarray.SIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          gpuarray.SSIZE,
          "int32",
          "int32",
      ]
      kernels.append(
          Kernel(
              code=code,
              name=kname,
              params=params,
              flags=flags,
              objvar=k_var,
          )
      )
      return kernels
  
imports: from __future__ import absolute_import, print_function, division
  from theano import Apply, Op
  from theano.tensor.extra_ops import CumOp
  from .basic_ops import as_gpuarray_variable, GpuKernelBase, Kernel, GpuReshape, infer_context_name, gpuarray_helper_inc_dir
  from .opt import register_opt, op_lifter, register_opt2
  from .type import gpu_context_type
  from theano.gof import ParamsType
  import theano.scalar
label_desc: Use image node_id 9 for calling the GpuCumOp obj's underlying member method code with example usage: obj.gpu_kernels(node, nodename) and returns: kernels