code: def _get_hist_kernel(dtype_str, nbins, offset):
      """
      Build a kernel to compute a 64 bin histogram.
  
      Use templating to generate a customized kernel depending on the input data type.
  
      Memoized to avoid compiling the same kernel twice.
      """
      type_str = _ew_types[dtype_str[1:]]
      from string import Template
  
      code = Template(
          _common_fp16_to_fp32
          + r"""
  
  #define MAX(a,b) (a > b ? a : b)
  #define MIN(a,b) (a < b ? a : b)
  
  __global__ void kernel_histo (
      int* d_hist, const $in_type* a1_in,
      int strides, int size)
  {
      const int tid = threadIdx.x;
      const int bid = blockIdx.x;
  
      __shared__ int s[$nbins];
      if(tid < $nbins){
          s[tid] = 0;
      }
  
      if(bid == 0 && tid < $nbins){
          d_hist[tid] = 0;
      }
  
      for (int i = tid + blockDim.x*bid; i < size; i += strides)
      {
          float a1 = $convert_to_float(__ldg(a1_in + i));
  
          float absval = fabs(a1);
  
          float logabs = round(log2f(absval));
  
          int bin = MIN($nbins-1, MAX(0, logabs-($offset)));
  
          atomicAdd(&s[bin], 1);
  
      }
  
      __syncthreads();
  
      if(tid < $nbins){
          atomicAdd(&d_hist[tid], s[tid]);
      }
  }
  """
      )
  
      module = SourceModule(
          code.substitute(
              in_type=type_str["type"],
              convert_to_float=type_str["cvt"],
              nbins=nbins,
              offset=offset,
          ),
          options=[],
      )
      kernel = module.get_function("kernel_histo")
      kernel.prepare("PPII")
      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 12 for calling a global function with example usage: _get_hist_kernel(dtype_str, nbins, offset) and returns: kernel