code: def _get_fprop_roipooling(clss):
      code = r"""
  #define FLT_MAX 3.402823466E+38F
  
  __global__ void fprop_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* bottom_data, const float* bottom_rois, float* top_data,
      int* argmax_data, const float spatial_scale) {
      for (int index = blockIdx.x * blockDim.x + threadIdx.x; \
          index < (nthreads); index += blockDim.x * gridDim.x){
          // (c, ph, pw, n) is an element in the pooled output
          int n = index % num_rois;
          int pw = (index / num_rois) % pooled_width;
          int ph = (index / num_rois / pooled_width) % pooled_height;
          int c = index / num_rois / pooled_width / pooled_height;
  
          bottom_rois += n * 5;
          int roi_batch_ind = bottom_rois[0];
          int roi_start_w = round(bottom_rois[1] * spatial_scale);
          int roi_start_h = round(bottom_rois[2] * spatial_scale);
          int roi_end_w = round(bottom_rois[3] * spatial_scale);
          int roi_end_h = round(bottom_rois[4] * spatial_scale);
  
          // 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 hstart = static_cast<int>(floor(static_cast<float>(ph)
                                              * bin_size_h));
          int wstart = static_cast<int>(floor(static_cast<float>(pw)
                                              * bin_size_w));
          int hend = static_cast<int>(ceil(static_cast<float>(ph + 1)
                                           * bin_size_h));
          int wend = static_cast<int>(ceil(static_cast<float>(pw + 1)
                                           * bin_size_w));
  
          // Add roi offsets and clip to input boundaries
          hstart = min(max(hstart + roi_start_h, 0), height);
          hend = min(max(hend + roi_start_h, 0), height);
          wstart = min(max(wstart + roi_start_w, 0), width);
          wend = min(max(wend + roi_start_w, 0), width);
          bool is_empty = (hend <= hstart) || (wend <= wstart);
  
          // Define an empty pooling region to be zero
          float maxval = is_empty ? 0 : -FLT_MAX;
          // If nothing is pooled, argmax = -1 causes nothing to be backprop'd
          int maxidx = -1;
  
          bottom_data += c * height * width * img_count;
  
          for (int h = hstart; h < hend; ++h) {
            for (int w = wstart; w < wend; ++w) {
              int bottom_index = h * width * img_count + w * img_count + roi_batch_ind;
              if (bottom_data[bottom_index] > maxval) {
                maxval = bottom_data[bottom_index];
                maxidx = bottom_index;
              }
            }
          }
          top_data[index] = maxval;
          argmax_data[index] = maxidx;
          // Notice the maxidx (from bottom_index) is relative to the dimension
          // (h, w, img_count) of the feature map, so max value is HWN
      }
  }
  
  """
  
      module = SourceModule(code)
      kernel = module.get_function("fprop_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 2 for calling a global function with example usage: _get_fprop_roipooling(clss) and returns: kernel