code: def _get_compensated_sum_kernel(dtype, rounding):
      _compensated_sum = r"""
  
  %(common)s
  
  __global__ void compensated_sum(unsigned* rand_state,
            %(type)s* a_sum,
            %(type)s* a_cmp,
      const %(type)s* a_add,
      float cmp_scale, float add_scale,
      int row_strd, int col_strd, int n, int mantissa_bits)
  {
      const int tid = threadIdx.x;
      const int bid = blockIdx.x;
  
      int offset = bid * row_strd + tid * col_strd;
      int inc    = 32 * col_strd;
  
      a_sum += offset;
      a_cmp += offset;
      a_add += offset;
  
      %(inits)s
  
      for (int i = tid; i < n; i += 32)
      {
          float s32 = %(cvt)s(__ldg((const %(type)s*)a_sum));
          float c32 = %(cvt)s(__ldg((const %(type)s*)a_cmp));
          float a32 = %(cvt)s(__ldg(a_add));
  
          // Adjust amount to add by previous compensation
          float y32 = a32 * add_scale - c32 * cmp_scale;
  
          // Do the accumulation and truncate to the storage type
          float rnd_sum = s32 + y32;
          %(rnd_sum)s
  
          // Convert accumulation back to fp32 so we can do more math on it
          float t32 = %(cvt)s(t16);
  
          // recover the low order bits that were lost in the truncation
          float rnd_cmp = (t32 - s32) - y32;
          %(rnd_cmp)s
  
          *a_sum = t16;
          *a_cmp = c16;
  
          a_sum += inc;
          a_cmp += inc;
          a_add += inc;
      }
      %(finish)s
  }
  """
      template_vals = dict()
      for key in ("common", "inits", "finish"):
          template_vals[key] = ""
  
      if dtype == "f2":
          template_vals["common"] += _common_fp16_to_fp32
  
      if rounding:
          template_vals["common"] += _common_urand_gen
          template_vals["common"] += _common_round["nearest"].get(
              dtype, ""
          )
          template_vals["inits"] += (
              _init_rand_func + _init_rand_round_func
          )
          template_vals["finish"] += _finish_rand_func
          mode = "random"
      else:
          mode = "nearest"
  
      template_vals["common"] += _common_round[mode].get(dtype, "")
  
      template_vals["type"] = _ew_types[dtype]["type"]
      template_vals["cvt"] = _ew_types[dtype]["cvt"]
  
      no_op = "float {0} = {1};"
  
      rnd_sum = _ew_strings["round"][mode].get(dtype, no_op)
      rnd_cmp = _ew_strings["round"]["nearest"].get(dtype, no_op)
  
      template_vals["rnd_sum"] = rnd_sum.format("t16", "rnd_sum")
      template_vals["rnd_cmp"] = rnd_cmp.format("c16", "rnd_cmp")
  
      code = _compensated_sum % template_vals
  
      # f = open("compensated_sum.cu", "w")
      # print >>f, code
      # f.close()
  
      module = SourceModule(code)
      kernel = module.get_function("compensated_sum")
      kernel.prepare("PPPPffiiii")
      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 10 for calling a global function with example usage: _get_compensated_sum_kernel(dtype, rounding) and returns: kernel