code: def _get_compound_kernel(type_args, compute_capability):
      """
      generate compound kernel for the optree from type_args
      """
  
      # from the stack, rebuild a mutable tree
      tree = _build_tree(type_args)
      # _print_tree(tree)
      # exit()
  
      # split all reductions and post reduction scalar operations out of the tree
      # sub-trees are converted to stacks and pushed onto stages list
      stages = _split_stages(tree)
      # _print_tree(tree)
      # exit()
  
      # set the final stage type to type of output (scalar or elementwise)
      last_stage = "red_out" if tree[1] == 1 else "ew_out"
      # convert the remainder of tree to stack
      stages.append((last_stage, _post_order(tree)))
  
      # for stage, stage_data in enumerate(stages):
      #     print stage_data[0], stage
      #     for s in stage_data[1]: print s
      #     print
      # exit()
  
      stack = list()
      placeholders = list()
      stage_out_reg = dict()
      arg_dict = dict()
      array_ids = set()
      fp16In = False
      rand_init = False
      rand_func = False
      threads = type_args[-1][3]
      template = _ew_template
      template_vals = {
          "threads": threads,
          "name": _get_kernel_name(),
          "common": list(),
          "inits": list(),
          "finish": list(),
      }
  
      for stage, stage_data in enumerate(stages):
          stage_type, stage_stack = stage_data
          new_placeholders = list()
  
          # build out the template as we process stages
          if stage_type == "reduction":
              new_placeholders.append("loads%d" % stage)
              new_placeholders.append("ops%d" % stage)
              new_placeholders.append("shfl_red%d" % stage)
              template += _stage_template["loop"].format(stage)
              if threads > 32:
                  new_placeholders.append("var_red%d" % stage)
                  new_placeholders.append("share1_red%d" % stage)
                  new_placeholders.append("share2_red%d" % stage)
                  template += _stage_template["red"].format(stage)
              else:
                  template += _stage_template["red32"].format(stage)
  
          elif stage_type == "scalar":
              new_placeholders.append("ops%d" % stage)
              template += _stage_template["red_ops"].format(stage)
  
          elif stage_type == "red_out":
              new_placeholders.append("ops%d" % stage)
              template += _stage_template["red_out"].format(stage)
  
          else:  # ew_out
              new_placeholders.append("loads%d" % stage)
              new_placeholders.append("ops%d" % stage)
              template += _stage_template["loop"].format(stage)
  
          for key in new_placeholders:
              template_vals[key] = []
          placeholders.extend(new_placeholders)
  
          for arg_i, arg in enumerate(stage_stack):
              arg_type, arg_id = arg[0:2]
  
              # Array operands
              if arg_type is ng.GPUTensor:
                  dtype, take_axis = arg[2:4]
  
                  is_out_tensor = (
                      True
                      if stage == len(stages) - 1 and arg_i == 0
                      else False
                  )
  
                  # first arg is output array, don't put on stack
                  if is_out_tensor:
                      out_dtype = dtype
                      out_take = take_axis
                  else:
                      stack.append("a%d" % arg_id)
  
                  # 0: arg_id, 1: stage, 2: type, 3: cvt
                  ew_dtype = _ew_types[dtype]
                  fmt = (
                      arg_id,
                      stage,
                      ew_dtype["type"],
                      ew_dtype["cvt"],
                  )
  
                  # First time we see a tensor initialize everything
                  if arg_id not in array_ids:
                      array_ids.add(arg_id)
                      array_ids.add((arg_id, stage))
  
                      sig = "Pii"
                      if take_axis > 0:
                          sig += "P"
  
                      # output tensor
                      if is_out_tensor:
                          ew_out = _ew_strings["out%d" % take_axis]
                          arguments = ew_out["arguments"].format(*fmt)
                          template_vals["inits"].append(
                              ew_out["inits"].format(*fmt)
                          )
                      # input tensors
                      else:
                          ew_in = _ew_strings["in%d" % take_axis]
                          loads = "loads%d" % stage
                          arguments = ew_in["arguments"].format(*fmt)
                          template_vals["inits"].append(
                              ew_in["inits"].format(*fmt)
                          )
                          template_vals[loads].append(
                              ew_in["loads"].format(*fmt)
                          )
  
                      if dtype == "f2" and not fp16In:
                          template_vals["common"].append(
                              _common_fp16_to_fp32
                          )
                          fp16In = True
  
                      arg_dict[arg] = (sig, arguments)
  
                  # Subsequent times we see a tensor just initialize inits and
                  # loads
                  elif (arg_id, stage) not in array_ids:
                      array_ids.add((arg_id, stage))
                      ew_in = _ew_strings["in%d" % take_axis]
                      loads = "loads%d" % stage
                      template_vals["inits"].append(
                          ew_in["inits"].format(*fmt)
                      )
                      template_vals[loads].append(
                          ew_in["loads"].format(*fmt)
                      )
  
              # Constant operands
              elif arg_type is float:
                  stack.append("c%d" % arg_id)
                  if arg not in arg_dict:
                      arg_dict[arg] = (
                          "f",
                          _ew_strings["const"]["arguments"].format(
                              arg_id
                          ),
                      )
  
              # Operations (arg_type = op_name)
              else:
                  if arg_type == "assign":
                      ops = "ops%d" % stage
  
                      # loop end condition for last stage
                      sig = "i"
                      arguments = ["const int n%d" % stage]
  
                      # rounding mode
                      if arg[2]:
                          mode = "random"
                          sig += "i"
                          arguments.append("const int mantissa_bits")
                          if not rand_init:
                              rand_init = _init_rand(template_vals)
                          template_vals["inits"].append(
                              _init_rand_round_func
                          )
                      else:
                          mode = "nearest"
  
                      arg_dict[arg] = (sig, ", ".join(arguments))
  
                      out_val = stack.pop()
                      # if the last stack value came from an argmax/min just do
                      # implicit type conversion
                      if out_val[0] == "i" and out_dtype[0] in "iu":
                          ew_round = None
                      else:
                          ew_round = _ew_strings["round"][mode].get(
                              out_dtype, None
                          )
                          ew_common = _common_round[mode].get(
                              out_dtype, None
                          )
                          if ew_common:
                              template_vals["common"].append(ew_common)
  
                      if ew_round:
                          round_val = "r%d" % arg_id
                          template_vals[ops].append(
                              ew_round.format(round_val, out_val)
                          )
                      else:
                          round_val = out_val
  
                      template_vals[ops].append(
                          _ew_strings["out%d" % out_take][
                              "output"
                          ].format(round_val)
                      )
  
                  elif arg in stage_out_reg:
                      stack.append(stage_out_reg[arg])
  
                  elif arg_type in _float_ops:
                      if len(template_vals["name"]) < 16:
                          template_vals["name"].append(arg_type)
  
                      ops = "ops%d" % stage
  
                      (num_ops, op_code) = _float_ops[arg_type]
  
                      if arg_type == "rand":
                          if not rand_init:
                              rand_init = _init_rand(template_vals)
                          if not rand_func:
                              template_vals["common"].append(
                                  _common_frand
                              )
                              rand_func = True
  
                      op_list = ["r%d" % arg_id]
  
                      # build the operands from the stack
                      for i in range(num_ops):
                          op_list.append(stack.pop())
  
                      if arg_type == "onehot":
                          hot_axis = arg[2]
                          test_val = "i" if hot_axis else "bid"
  
                          ew_in = _ew_strings[
                              arg_type + native_str(hot_axis)
                          ]
                          loads = "loads%d" % stage
                          template_vals["inits"].append(
                              ew_in["inits"].format(arg_id)
                          )
                          template_vals[loads].append(
                              ew_in["loads"].format(arg_id)
                          )
                          op_list.append("onehot%d" % arg_id)
                          op_list.append(test_val)
  
                          arg_dict[arg] = (
                              "P",
                              ew_in["arguments"].format(arg_id),
                          )
  
                      template_vals[ops].append(
                          op_code.format(*op_list)
                      )
  
                      # if this is the last op on the current stack, store its register stage
                      # in the stage output dict
                      if arg_i == len(stage_stack) - 1:
                          stage_out_reg[arg] = op_list[0]
                      # otherwise push the reg onto the stack as normal
                      else:
                          stack.append(op_list[0])
  
                  elif arg_type in _reduction_ops:
                      if len(template_vals["name"]) < 16:
                          template_vals["name"].append(arg_type)
  
                      # loop end condition for current stage
                      # add regardless of duplicate reduction stage
                      arg_dict[arg] = ("i", "const int n%d" % stage)
  
                      # avoid float conversion for argmax/min
                      reg = "i" if "arg" == arg_type[0:3] else "r"
  
                      ops = "ops%d" % stage
                      shfl_red = "shfl_red%d" % stage
                      red_arg = "%s%d" % (reg, arg_id)
                      red_strings = _reduction_ops[arg_type]
                      stack_arg = stack.pop()
  
                      template_vals["inits"].append(
                          red_strings["inits"].format(red_arg)
                      )
                      template_vals[ops].append(
                          red_strings["ops"].format(red_arg, stack_arg)
                      )
                      template_vals[shfl_red].append(
                          red_strings["shfl_red"].format(red_arg)
                      )
                      if threads > 32:
                          var_red = "var_red%d" % stage
                          shr1_red = "share1_red%d" % stage
                          shr2_red = "share2_red%d" % stage
                          template_vals[var_red].append(red_arg)
                          template_vals[shr1_red].append(
                              red_strings["share1_red"].format(red_arg)
                          )
                          template_vals[shr2_red].append(
                              red_strings["share2_red"].format(red_arg)
                          )
  
                      # reduction ops are always the last on the stack
                      # just store the register state in the stage output dict
                      stage_out_reg[arg] = red_arg
  
                  else:
                      raise ValueError("Bad op type.")
  
      if (
          compute_capability[0] == 3 and compute_capability[1] < 5
      ) or compute_capability[0] < 3:
          template_vals["common"].append(_common_kepler)
  
      template += _fin_template
  
      # since we reorderd the operations we need to generate the argument list
      # in the original order
      sig = "P"
      arguments = list()
      unused = 1
      for arg in type_args:
          params = arg_dict.get(arg, False)
          if params:
              sig += params[0]
              arguments.append(params[1])
              del arg_dict[arg]
          # fill in the loop counter for the duplicate reductions that were
          # removed
          elif arg[0] in _reduction_ops:
              sig += "i"
              arguments.append("const int unused%d" % unused)
              unused += 1
  
      # convert lists to strings
      template_vals["name"] = "_".join(template_vals["name"])
      template_vals["common"] = "\n".join(template_vals["common"])
      template_vals["arguments"] = ",\n    ".join(arguments)
      template_vals["inits"] = "\n    ".join(template_vals["inits"])
      template_vals["finish"] = "\n".join(template_vals["finish"])
  
      # add the dynamic placeholders: loads#, ops#, reduction#
      for key in placeholders:
          template_vals[key] = "\n        ".join(template_vals[key])
  
      # populate the template
      code = template % template_vals
  
      # debugging:
      # print "Compiling %s" % template_vals["name"]
      # f = open("kernel.cu", "w")
      # f = open("%s.cu" % template_vals["name"], "w")
      # print >>f, code
      # f.close()
  
      # ,"-G" , keep=False
      # module = SourceModule(code, options=["--use_fast_math"])
      module = SourceModule(code, options=[])
      kernel = module.get_function(template_vals["name"])
      kernel.name = template_vals["name"]
      kernel.prepare(sig)
  
      return kernel
  
imports: from __future__ import division
  from future.utils import native_str
  import os.path
  import re
  import traceback
  import numpy
  from pycuda.tools import context_dependent_memoize
  from pytools import memoize
  from neon import logger
  import neon.backends.nervanagpu
  from neon.backends.util.source_module import SourceModule
  from neon.backends.cuda_templates import _ew_template, _stage_template, _fin_template, _init_rand_func, _init_rand_round_func, _finish_rand_func, _common_kepler, _common_urand_gen, _common_frand, _common_round, _common_fp16_to_fp32, _ew_types, _ew_strings, _is_finite, _float_ops, _reduction_ops
  from neon.backends.cuda_batchnorm import _get_bn_fprop_kernel, _get_bn_bprop_kernel
  from neon.backends.kernels.cuda.lookuptable import _get_lut_bprop_kernel, _get_sorting_kernel
label_desc: Use image node_id 7 for calling a global function with example usage: _get_compound_kernel(type_args, compute_capability) and returns: kernel