code: def _get_bprop_avg(clss, compute_capability):
      code = r"""
  
  %(common)s
  
  __global__ void spool_bprop_avg(
      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 P, int Q,
      int magic_P, int shift_P, int QN, int PQN, int MPQN,
      int pad_c, int pad_d, int pad_h, int pad_w,
      int str_c, int str_d, int str_h, int str_w,
      int S, 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 supP, int supQ, int shlP, int maskP, int shrP,
      int shlQ, int maskQ, int shrQ, int maskN, int shrN
      %(stats_args)s
      )
  {
      __shared__ float rcpWindowSize[32];
      extern __shared__ int lut[];
  
      int tid = threadIdx.x;
  
      int q  = blockIdx.x;
      int mp = blockIdx.y;
      int k  = blockIdx.z;
  
      int m = mp * magic_P; m >>= shift_P;
      int p = mp - m*supP;
  
      // zigzag q back and forth to improve L2 cache perf
      if (p & 1)
          q = supQ - q - 1;
  
      // Superblock P and Q
      p = (p << shlP) + ((tid & maskP) >> shrP);
      q = (q << shlQ) + ((tid & maskQ) >> shrQ);
      int n = tid & maskN;
  
      int sb = tid >> shrN;
  
      O += n;
      I += k*MPQN + m*PQN + p*QN + mad16(q, N, n);
  
      float delta = 0.0f;
      if (p < P && q < Q && n < N)
          delta  = %(cvt)s(__ldg(I));
  
      if (tid < 32)
      {
          int kj = k * str_c - pad_c;
          int mt = m * str_d - pad_d;
          int pr = p * str_h - pad_h;
          int qs = q * str_w - pad_w;
  
          int inc    = min(maskN + 1, 32);
          int sbBits = 1 << min(shrN, 5);
          int sbMask = ~(-1 << sbBits) << mad16(sb, sbBits, 0);
  
          int window_size = 0;
          int jrst = n;
          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 x = qs + s;
              int y = pr + r;
              int z = mt + t;
              int c = kj + j;
  
              bool bounds_x  = x >= 0 && x < W;
              bool bounds_y  = y >= 0 && y < H;
              bool bounds_z  = z >= 0 && z < D;
              bool bounds_c  = c >= 0 && c < C;
              bool in_bounds = bounds_x && bounds_y && bounds_z && bounds_c;
  
              // Count the total valid slices
              window_size += __popc(sbMask & __ballot(in_bounds));
  
              int sliceI  = c*DHWN + z*HWN + y*WN + x*N;
  
              int lut_offset = mad16(sb, JRST, jrst);
  
              lut[lut_offset] = in_bounds ? sliceI : -1;
              jrst += inc;
          }
          // TODO confirm kepler OK
          unsigned int shrN_mask = (shrN < 32) ? max(0, ((1 << shrN) - 1)) : 0xffffffff;
          if((tid & shrN_mask) == 0)
              rcpWindowSize[sb] = 1.0f / (float)window_size;
      }
      __syncthreads();
  
      int intermediate_max = 0;
  
      if (p < P && q < Q && n < N)
      {
          delta *= alpha * rcpWindowSize[sb];
          bool load_beta = beta != 0.0f;
          int jrst = 0;
  
          while (jrst < JRST)
          {
              int lut_offset = mad16(sb, JRST, jrst);
  
              int offset0 = lut[lut_offset + 0];
              int offset1 = lut[lut_offset + 1];
              int offset2 = lut[lut_offset + 2];
              int offset3 = lut[lut_offset + 3];
  
              %(type)s* out0 = O + offset0;
              %(type)s* out1 = O + offset1;
              %(type)s* out2 = O + offset2;
              %(type)s* out3 = O + offset3;
  
              bool valid0 = jrst + 0 < JRST && offset0 >= 0;
              bool valid1 = jrst + 1 < JRST && offset1 >= 0;
              bool valid2 = jrst + 2 < JRST && offset2 >= 0;
              bool valid3 = jrst + 3 < JRST && offset3 >= 0;
  
              float beta0 = valid0 && load_beta ? %(cvt)s(__ldg(out0)) * beta : 0.0f;
              float beta1 = valid1 && load_beta ? %(cvt)s(__ldg(out1)) * beta : 0.0f;
              float beta2 = valid2 && load_beta ? %(cvt)s(__ldg(out2)) * beta : 0.0f;
              float beta3 = valid3 && load_beta ? %(cvt)s(__ldg(out3)) * beta : 0.0f;
  
              %(type)s temp_out0 = valid0 ? %(cvt_out)s(%(mul_by_scale)s(delta + beta0)) : 0.0f;
              %(type)s temp_out1 = valid1 ? %(cvt_out)s(%(mul_by_scale)s(delta + beta1)) : 0.0f;
              %(type)s temp_out2 = valid2 ? %(cvt_out)s(%(mul_by_scale)s(delta + beta2)) : 0.0f;
              %(type)s temp_out3 = valid3 ? %(cvt_out)s(%(mul_by_scale)s(delta + beta3)) : 0.0f;
  
              // predicate writes with no-op flag.
              if (!(flags & 1)) {
                  if (valid0) *out0 = temp_out0;
                  if (valid1) *out1 = temp_out1;
                  if (valid2) *out2 = temp_out2;
                  if (valid3) *out3 = temp_out3;
              }
              intermediate_max = max_abs(intermediate_max, temp_out0);
              intermediate_max = max_abs(intermediate_max, temp_out1);
              intermediate_max = max_abs(intermediate_max, temp_out2);
              intermediate_max = max_abs(intermediate_max, temp_out3);
  
              jrst += 4;
          }
      }
      intermediate_max += 0;
      %(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")
      kernel.prepare("3P 2f 44I" + ("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 8 for calling a global function with example usage: _get_bprop_avg(clss, compute_capability) and returns: kernel