code: def _get_bprop_avg_overlap_smallN(clss, compute_capability):
      code = r"""
  
  %(common)s
  
  union LutEntry {
      struct {
          int   slice;
          float rcp_in;
      } data;
      int2 data2;
  };
  
  __device__ __forceinline__ int imin(int val1, int val2)
  {
      int ret;
      asm("min.s32 %%0, %%1, %%2;" : "=r"(ret) : "r"(val1), "r"(val2));
      return ret;
  }
  
  __global__ void spool_bprop_avg_overlap_smallN(
      const %(type)s* I, %(type)s* O, const unsigned char* A,
      float alpha, float beta, 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,
      int supH, int supW, int shlH, int maskH, int shrH,
      int shlW, int maskW, int shrW, int maskN, int shrN, int maxLutSize
      %(stats_args)s  // template for "int* maxabs, float scale0"
      )
  {
      extern __shared__ int2 lut[];
  
      int tid = threadIdx.x;
  
      int x  = blockIdx.x;
      int yz = blockIdx.y;
      int c  = blockIdx.z;
  
      int z = yz * magic_H; z >>= shift_H;
      int y = yz - z*supH;
  
      // zigzag q back and forth to improve L2 cache perf
      if (y & 1)
          x = supW - x - 1;
  
      // Superblock H and W
      y = (y << shlH) + ((tid & maskH) >> shrH);
      x = (x << shlW) + ((tid & maskW) >> shrW);
      int n = tid & maskN;
  
      int sb = tid >> shrN;
  
      I += n;
      O += c*DHWN + z*HWN + y*WN + mad16(x, N, n);
  
      float O_val = beta != 0.0f && y < H && x < W && n < N ? %(cvt)s(__ldg(O)) : 0.0f;
  
      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 sbSize = maskN + 1;
      int sbBits = mad16(sb, sbSize, 0);
      unsigned sbMask = ~(0xffffffff << sbSize) << sbBits;
      unsigned dep_thd_mask = (0xffffffff >> (32 - n)) << sbBits;
  
      int lut_offset = mad16(sb, maxLutSize, 0);
      int lut_size = 0;
      int jrst = n;
      int JRSTend = JRST;
      if (JRSTend & maskN)
          JRSTend += sbSize - (JRSTend & maskN);
  
      while (jrst < JRSTend)
      {
          int j   = div16(jrst, magic_RST, shift_RST);
          int rst = mod16(jrst, j, RST);
  
          int t   = div16(rst, magic_RS, shift_RS);
          int rs  = mod16(rst, t, RS);
  
          int r   = div16(rs, magic_S, shift_S);
          int s   = mod16(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        = div16(k_prime, magic_str_c, shift_str_c);
          int  k_mod    = mod16(k_prime, k, str_c);
          bool k_bounds = k_mod == 0 && k >= 0 && k < K;
  
          int  m        = div16(m_prime, magic_str_d, shift_str_d);
          int  m_mod    = mod16(m_prime, m, str_d);
          bool m_bounds = m_mod == 0 && m >= 0 && m < M;
  
          int  p        = div16(p_prime, magic_str_h, shift_str_h);
          int  p_mod    = mod16(p_prime, p, str_h);
          bool p_bounds = p_mod == 0 && p >= 0 && p < P;
  
          int  q        = div16(q_prime, magic_str_w, shift_str_w);
          int  q_mod    = mod16(q_prime, q, str_w);
          bool q_bounds = q_mod == 0 && q >= 0 && q < Q;
  
          bool in_bounds = jrst < JRST && k_bounds && m_bounds && p_bounds && q_bounds;
  
          // Get a mask of all valid slices in the warp
          unsigned ballot = __ballot(in_bounds);
  
          // Count the total valid slices in this superblock
          int sb_slices = __popc(sbMask & ballot);
  
          if (in_bounds)
          {
              // Count all the valid slices below this threadid
              int dep_thd_cnt = __popc(dep_thd_mask & ballot);
  
              int c_left = msub16(k, str_c, pad_c);
              int z_left = msub16(m, str_d, pad_d);
              int y_left = msub16(p, str_h, pad_h);
              int x_left = msub16(q, str_w, pad_w);
  
              float k_in = (float)imin( imin(J + c_left, J), imin(C - c_left, J) );
              float m_in = (float)imin( imin(T + z_left, T), imin(D - z_left, T) );
              float p_in = (float)imin( imin(R + y_left, R), imin(H - y_left, R) );
              float q_in = (float)imin( imin(S + x_left, S), imin(W - x_left, S) );
  
              float total_in = q_in * p_in * m_in * k_in;
  
              LutEntry entry;
              entry.data.slice  = k*MPQN + m*PQN + p*QN + mad16(q, N, 0);
              entry.data.rcp_in = total_in > 0.0f ? 1.0f / total_in : 0.0f;
  
              lut[lut_offset + lut_size + dep_thd_cnt] = entry.data2;
          }
          lut_size += sb_slices;
          jrst += sbSize;
      }
  
      int intermediate_max = 0;
  
      if (y < H && x < W && n < N)
      {
          int jrst = 0;
          float out = 0.0f;
  
          while (jrst < maxLutSize)
          {
              LutEntry entry0;
              LutEntry entry1;
              LutEntry entry2;
              LutEntry entry3;
  
              lut_offset = mad16(sb, maxLutSize, jrst);
  
              entry0.data2 = lut[lut_offset + 0];
              entry1.data2 = lut[lut_offset + 1];
              entry2.data2 = lut[lut_offset + 2];
              entry3.data2 = lut[lut_offset + 3];
  
              out += jrst + 0 < lut_size ? %(cvt)s(__ldg(I + entry0.data.slice)) * entry0.data.rcp_in : 0.0f;
              out += jrst + 1 < lut_size ? %(cvt)s(__ldg(I + entry1.data.slice)) * entry1.data.rcp_in : 0.0f;
              out += jrst + 2 < lut_size ? %(cvt)s(__ldg(I + entry2.data.slice)) * entry2.data.rcp_in : 0.0f;
              out += jrst + 3 < lut_size ? %(cvt)s(__ldg(I + entry3.data.slice)) * entry3.data.rcp_in : 0.0f;
  
              jrst += 4;
          }
          %(type)s temp_out = %(cvt_out)s( %(mul_by_scale)s (out*alpha + O_val*beta));
          if (!(flags & 1))
              *O = temp_out;
  
          // max-abs over unrolls
          intermediate_max = max_abs(intermediate_max, temp_out);  // used for abs
      }
      intermediate_max += 0;
      %(atomic_max)s
  
  }
  """
      template_vals = prepare_template_vals(clss, compute_capability)
      code = code % template_vals
      # f = open("pool.cu", "w")
      # print >>f, code
      # f.close()
      module = SourceModule(code, options=["--use_fast_math"])
      kernel = module.get_function("spool_bprop_avg_overlap_smallN")
      kernel.prepare("3P 2f 58I" + ("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 12 for calling a global function with example usage: _get_bprop_avg_overlap_smallN(clss, compute_capability) and returns: kernel