code: def _get_nms_kernel():
      code = r"""
  #define DIVUP(m,n) ((m) / (n) + ((m) % (n) > 0))
  int const threadsPerBlock = sizeof(unsigned int) * 8;
  
  __device__ inline float devIoU(float const * const a, float const * const b,
                                 int const offset) {
    float left = max(a[0], b[0]), right = min(a[2], b[2]);
    float top = max(a[1], b[1]), bottom = min(a[3], b[3]);
    float width = max(right - left + offset, 0.f), height = max(bottom - top + offset, 0.f);
    float interS = width * height;
    float Sa = (a[2] - a[0] + offset) * (a[3] - a[1] + offset);
    float Sb = (b[2] - b[0] + offset) * (b[3] - b[1] + offset);
  
    return interS / (Sa + Sb - interS);
  }
  
  __global__ void nms_kernel(const int n_boxes, const float nms_overlap_thresh,
                             const float *dev_boxes, unsigned int *dev_mask,
                             const bool normalized) {
    const int row_start = blockIdx.y;
    const int col_start = blockIdx.x;
  
    // if (row_start > col_start) return;
  
    const int row_size =
          min(n_boxes - row_start * threadsPerBlock, threadsPerBlock);
    const int col_size =
          min(n_boxes - col_start * threadsPerBlock, threadsPerBlock);
  
    // if boxes are not normalized to image dim, we use an offset of 1 in
    // calculating the box width and height.
    const float offset = normalized ? 0 : 1;
  
    __shared__ float block_boxes[threadsPerBlock * 5];
    if (threadIdx.x < col_size) {
      block_boxes[threadIdx.x * 5 + 0] =
          dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 0];
      block_boxes[threadIdx.x * 5 + 1] =
          dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 1];
      block_boxes[threadIdx.x * 5 + 2] =
          dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 2];
      block_boxes[threadIdx.x * 5 + 3] =
          dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 3];
      block_boxes[threadIdx.x * 5 + 4] =
          dev_boxes[(threadsPerBlock * col_start + threadIdx.x) * 5 + 4];
    }
    __syncthreads();
  
    if (threadIdx.x < row_size) {
      const int cur_box_idx = threadsPerBlock * row_start + threadIdx.x;
      const float *cur_box = dev_boxes + cur_box_idx * 5;
      int i = 0;
      unsigned int t = 0;
      int start = 0;
      if (row_start == col_start) {
        start = threadIdx.x + 1;
      }
      for (i = start; i < col_size; i++) {
        if (devIoU(cur_box, block_boxes + i * 5, offset) > nms_overlap_thresh) {
          t |= 1UL << i;
        }
      }
      const int col_blocks = DIVUP(n_boxes, threadsPerBlock);
      dev_mask[cur_box_idx * col_blocks + col_start] = t;
    }
  }
  
  """
  
      module = SourceModule(code)
      kernel = module.get_function("nms_kernel")
      sig = "1I 1f 2P 1b"
      kernel.prepare(sig)
      return kernel
  
label_desc: Use image node_id 1 for calling a global function with example usage: _get_nms_kernel() and returns: kernel