code: def _get_bprop_avg_overlap(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(
      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
      %(stats_args)s  // template for "int* maxabs, float scale0"
      )
  {
      int    __shared__ lutSize;
      extern __shared__ int2 lut[];
  
      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;
  
      I += n;
      O += c*DHWN + z*HWN + y*WN + x*N + n;
  
      float O_val = beta != 0.0f ? %(cvt)s(__ldg(O)) : 0.0f;
      int lut_size;
  
      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;
  
          unsigned dep_thd_mask = 0xffffffff;
          dep_thd_mask >>= 32 - tid;
  
          lut_size = 0;
  
          int jrst = tid;
          while (jrst < JRST)
          {
              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 = 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
              unsigned warp_slices = __popc(ballot);
  
              if (in_bounds)
              {
                  // Count all the valid slices below this threadid
                  unsigned 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;
  
                  float rcp_in = total_in > 0.0f ? 1.0f / total_in : 0.0f;
  
                  LutEntry entry;
                  entry.data.slice  = k*MPQN + m*PQN + p*QN + mad16(q, N, 0);
                  entry.data.rcp_in = rcp_in;
  
                  lut[lut_size + dep_thd_cnt] = entry.data2;
              }
              lut_size += warp_slices;
              jrst += 32;
          }
          if(tid==0)
              lutSize = lut_size;
      }
      __syncthreads();
  
      lut_size = lutSize;
  
      int jrst = 0;
      float out = 0.0f;
      int intermediate_max = 0;
  
      while (jrst < lut_size)
      {
          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];
  
          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
  
      %(atomic_max)s
  
  }
  """
      template_vals = prepare_template_vals(clss, compute_capability)
      code = code % template_vals
      module = SourceModule(code, options=["--use_fast_math"])
      kernel = module.get_function("spool_bprop_avg_overlap")
      kernel.prepare("3P 2f 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 10 for calling a global function with example usage: _get_bprop_avg_overlap(clss, compute_capability) and returns: kernel