

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

dataset/ETHPy150Open NervanaSystems/neon/neon/backends/kernels/cuda/roipooling.py _get_fprop_roipooling/VarMisuse@37/1455 `kernel`->`sig`