code: def _get_bprop_roipooling(clss):
      code = r"""
  __global__ void bprop_roipooling(const int nthreads,
      const int num_rois, const int img_count,
      const int channels, const int height, const int width,
      const int pooled_height, const int pooled_width,
      const float* top_diff, const float* bottom_rois, float* bottom_diff,
      const int* argmax_data, const float spatial_scale) {
      for (int index = blockIdx.x * blockDim.x + threadIdx.x; \
          index < (nthreads); index += blockDim.x * gridDim.x){
          // (c, h, w, n) coords in bottom data on feature map
          int n = index % img_count;
          int w = (index / img_count) % width;
          int h = (index / img_count / width) % height;
          int c = index / img_count/ width / height;
  
          float gradient = 0;
          // Accumulate gradient over all ROIs that pooled this element
          for (int roi_n = 0; roi_n < num_rois; ++roi_n) {
            const float* offset_bottom_rois = bottom_rois + roi_n * 5;
            int roi_batch_ind = offset_bottom_rois[0];
            // Skip if ROI's batch index doesn't match n
            if (n != roi_batch_ind) {
              continue;
            }
  
            int roi_start_w = round(offset_bottom_rois[1] * spatial_scale);
            int roi_start_h = round(offset_bottom_rois[2] * spatial_scale);
            int roi_end_w = round(offset_bottom_rois[3] * spatial_scale);
            int roi_end_h = round(offset_bottom_rois[4] * spatial_scale);
  
            // Skip if ROI doesn't include (h, w)
            const bool in_roi = (w >= roi_start_w && w <= roi_end_w &&
                                 h >= roi_start_h && h <= roi_end_h);
            if (!in_roi) {
              continue;
            }
  
            int offset = c * pooled_height * pooled_width * num_rois;
            const float* offset_top_diff = top_diff + offset;
            const int* offset_argmax_data = argmax_data + offset;
  
            // Compute feasible set of pooled units that could have pooled
            // this bottom unit
  
            // Force malformed ROIs to be 1x1
            int roi_width = max(roi_end_w - roi_start_w + 1, 1);
            int roi_height = max(roi_end_h - roi_start_h + 1, 1);
  
            float bin_size_h = static_cast<float>(roi_height)
                               / static_cast<float>(pooled_height);
            float bin_size_w = static_cast<float>(roi_width)
                               / static_cast<float>(pooled_width);
  
            int phstart = floor(static_cast<float>(h - roi_start_h) / bin_size_h);
            int phend = ceil(static_cast<float>(h - roi_start_h + 1) / bin_size_h);
            int pwstart = floor(static_cast<float>(w - roi_start_w) / bin_size_w);
            int pwend = ceil(static_cast<float>(w - roi_start_w + 1) / bin_size_w);
  
            phstart = min(max(phstart, 0), pooled_height);
            phend = min(max(phend, 0), pooled_height);
            pwstart = min(max(pwstart, 0), pooled_width);
            pwend = min(max(pwend, 0), pooled_width);
  
            for (int ph = phstart; ph < phend; ++ph) {
              for (int pw = pwstart; pw < pwend; ++pw) {
                int top_index = ph * pooled_width * num_rois + pw * num_rois + roi_n;
                int bottom_index = h * width * img_count + w * img_count + roi_batch_ind;
                if (offset_argmax_data[top_index] == bottom_index) {
                  gradient += offset_top_diff[top_index];
                }
              }
            }
          }
          bottom_diff[index] = gradient;
      }
  }
  
  """
  
      module = SourceModule(code)
      kernel = module.get_function("bprop_roipooling")
      sig = "8I 4P 1f"
      kernel.prepare(sig)
      return kernel
  
imports: from neon.backends.util.source_module import SourceModule
  from pycuda.tools import context_dependent_memoize
label_desc: Use image node_id 3 for calling a global function with example usage: _get_bprop_roipooling(clss) and returns: kernel