code: def _get_bprop_lrn_overlap(clss, compute_capability):
      code = r"""
  %(common)s
  
  union LutEntry {
      struct {
          int slice;
          int argmax;
      } data;
      int2 data2;
  };
  
  __global__ void spool_bprop_lrn_overlap(
      const %(type)s* I, const %(type)s* O, const %(type)s* E, %(type)s* delta, const %(type)s* A,
      float alpha, float beta, float ascale, float bpower, int flags,
      int N, int W, int H, int D, int C,
      int WN, int HWN, int DHWN,
      int magic_H, int shift_H,
      int pad_w, int pad_h, int pad_d, int pad_c,
      int str_w, int str_h, int str_d, int str_c,
      int magic_str_w, int shift_str_w,
      int magic_str_h, int shift_str_h,
      int magic_str_d, int shift_str_d,
      int magic_str_c, int shift_str_c,
      int S, int R, int T, int J, int RS, int RST, int JRST,
      int magic_S, int shift_S, int magic_RS, int shift_RS,
      int magic_RST, int shift_RST,
      int Q, int P, int M, int K, int QN, int PQN, int MPQN
      %(stats_args)s  // template for "int* maxabs, float scale0"
      )
  {
      extern __shared__ int2 lut[];
      __shared__ float rcpWindowSize;
  
      int tid = threadIdx.x;
  
      int n  = tid;
      int x  = blockIdx.x;
      int yz = blockIdx.y;
      int c  = blockIdx.z;
  
      int z = yz * magic_H; z >>= shift_H;
      int y = yz - z*H;
  
      // zigzag q back and forth to improve L2 cache perf
      if (y & 1)
          x = W - x - 1;
  
      // O E A used inside JRST loop
      O += n;  // output
      E += n;  // error
      A += n;  // denom
  
      // I E A used for output
      const %(type)s* E_out = E;
      const %(type)s* A_out = A;
      delta += c*DHWN + z*HWN + y*WN + x*N + n;
      I     += c*DHWN + z*HWN + y*WN + x*N + n;
      E_out += c*DHWN + z*HWN + y*WN + x*N;
      A_out += c*DHWN + z*HWN + y*WN + x*N;
  
      float delta_val = (beta != 0.0f) ? %(cvt)s(__ldg(delta)) : 0.0f;
  
      // build the lookup table
      if (tid < 32)
      {
          int kj = c - J + pad_c + 1;
          int mt = z - T + pad_d + 1;
          int pr = y - R + pad_h + 1;
          int qs = x - S + pad_w + 1;
  
          int jrst = tid;
          while (jrst < JRST)
          {
              int j = jrst * magic_RST; j >>= shift_RST;
              int rst = jrst - j * RST;
  
              int t = rst * magic_RS; t >>= shift_RS;
              int rs = rst - t * RS;
  
              int r = rs * magic_S; r >>= shift_S;
              int s = rs - r*S;
  
              int k_prime = kj + j;
              int m_prime = mt + t;
              int p_prime = pr + r;
              int q_prime = qs + s;
  
              int k     = k_prime * magic_str_c; k >>= shift_str_c;
              int k_mod = k_prime - k*str_c;
              bool k_bounds = k_mod == 0 && k >= 0 && k < K;
  
              int m     = m_prime * magic_str_d; m >>= shift_str_d;
              int m_mod = m_prime - m*str_d;
              bool m_bounds = m_mod == 0 && m >= 0 && m < M;
  
              int p     = p_prime * magic_str_h; p >>= shift_str_h;
              int p_mod = p_prime - p*str_h;
              bool p_bounds = p_mod == 0 && p >= 0 && p < P;
  
              int q     = q_prime * magic_str_w; q >>= shift_str_w;
              int q_mod = q_prime - q*str_w;
              bool q_bounds = q_mod == 0 && q >= 0 && q < Q;
  
              bool in_bounds = k_bounds && m_bounds && p_bounds && q_bounds;
  
              int j_prime = c - k_prime + pad_c;
              int t_prime = z - m_prime + pad_d;
              int r_prime = y - p_prime + pad_h;
              int s_prime = x - q_prime + pad_w;
  
              int sliceI  = k*MPQN + m*PQN + p*QN + q*N;
              int argmaxI = j_prime*RST + t_prime*RS + r_prime*S + s_prime;
  
              LutEntry entry;
              entry.data.slice  = sliceI;
              entry.data.argmax = in_bounds ? argmaxI : -1;
  
              lut[jrst] = entry.data2;
              jrst += 32;
          }
  
          if(tid == 0)
          {
              rcpWindowSize = (float)RST/(float)JRST;
          }
      }
      __syncthreads();
  
      int jrst = 0;
      // float out = 0.0f;
      float array_delta = 0.0f;
      int intermediate_max = 0;
  
      while (jrst < JRST)
      {
          LutEntry entry0;
          LutEntry entry1;
          LutEntry entry2;
          LutEntry entry3;
  
          entry0.data2 = lut[jrst + 0];
          entry1.data2 = lut[jrst + 1];
          entry2.data2 = lut[jrst + 2];
          entry3.data2 = lut[jrst + 3];
  
          array_delta += (jrst + 0 < JRST && entry0.data.argmax >= 0) ? %(cvt)s(__ldg(O + entry0.data.slice)) * %(cvt)s(__ldg(E + entry0.data.slice)) / %(cvt)s(__ldg(A + entry0.data.slice)) : 0.0f;
          array_delta += (jrst + 1 < JRST && entry1.data.argmax >= 0) ? %(cvt)s(__ldg(O + entry1.data.slice)) * %(cvt)s(__ldg(E + entry1.data.slice)) / %(cvt)s(__ldg(A + entry1.data.slice)) : 0.0f;
          array_delta += (jrst + 2 < JRST && entry2.data.argmax >= 0) ? %(cvt)s(__ldg(O + entry2.data.slice)) * %(cvt)s(__ldg(E + entry2.data.slice)) / %(cvt)s(__ldg(A + entry2.data.slice)) : 0.0f;
          array_delta += (jrst + 3 < JRST && entry3.data.argmax >= 0) ? %(cvt)s(__ldg(O + entry3.data.slice)) * %(cvt)s(__ldg(E + entry3.data.slice)) / %(cvt)s(__ldg(A + entry3.data.slice)) : 0.0f;
  
          jrst += 4;
      }
  
      array_delta = -2 * bpower * ascale * __ldg(I) * array_delta * rcpWindowSize + (__ldg(E_out) * powf(__ldg(A_out), -bpower));
  
      %(type)s temp_out = %(cvt_out)s( %(mul_by_scale)s (array_delta*alpha + delta_val*beta));
      if (!(flags & 1)) {
          *delta = temp_out;
      }
  
      // compute max-abs
      intermediate_max = max_abs(intermediate_max, temp_out);  // used for abs
      %(atomic_max)s
  }
  """
  
      template_vals = prepare_template_vals(clss, compute_capability)
      code = code % template_vals
      module = SourceModule(code)
      kernel = module.get_function("spool_bprop_lrn_overlap")
      kernel.prepare("5P 4f 47I" + ("Pf" if (clss[0] == "x") else ""))
      return kernel
  
file_path: neon/neon/backends/kernels/cuda/pooling.py
imports: from builtins import str
  from pycuda.tools import context_dependent_memoize
  from neon.backends import cuda_templates
  from neon.backends.cuda_templates import _common_fp16_to_fp32, _common_round, _common_max_abs, _common_kepler, _ew_types
  from neon.backends.util.source_module import SourceModule
label_desc: Use image node_id 6 for calling a global function with example usage: _get_bprop_lrn_overlap(clss, compute_capability) and returns: kernel