

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

dataset/ETHPy150Open NervanaSystems/neon/neon/backends/kernels/cuda/roipooling.py _get_bprop_roipooling/VarMisuse@21/1507 `code`->`module`