[TOPI] Parallelize GPU NMS inner loop

[masahi]

This is a follow-up to #7136. I found a simple way to parallelize the inner loop of GPU NMS, which is currently done in a sequential way since #6839 and hence extremely slow if the number of input box is large. This change brings massive speedup on object detection models from PyTorch and Gluon, as shown below.

GPU NMS workload from PyTorch MaskRCNN (4500 boxes)
Before: 2.1 sec
After: 5.78 milli sec 

Workload from Gluon SSD
Before: 12.5 milli sec
After: 0.206 milli sec

Before I explain what I did, here is how currently we do the sequential, O(N ** 2) triangle loop. This is done by a single thread.

# The outer loop goes through boxes sorted by their score.
for j in range(nboxes):
    # The inner loop check if box j does not overlap with all other preceding boxes k
    for k in range(j):
        if box k is valid:
            do IOU test between j and k and possibly invalidate j

My parallelization instead does the above triangle in the following way:

for j in range(nboxes):
   if j is still valid:
        for k in range(j + 1, nboxes) in parallel:
             do IOU test between j and k and possibly invalidate box k
   # Flush IOU test results from the inner loop to global memory before continuing to the next loop
   # So that other threads can also see them
   syncthreads() 

The idea is, at the start of the inner loop, box j is assumed to be a valid box, and the inner loop invalidates other succeeding boxes that have high overlap with the newly found valid box j. The inner loop can be trivially done in parallel and the number of IOU tests reduces to O(# selected boxes * N).

Now, the inner loop is done in parallel and the other loop is sequential. All threads need to do the outer loop in a lock step: The results of checking if box j is still valid must be consistent across all threads. Since we cannot do global sync inside kernels, I use only one thread block for parallelization and use __syncthreads() after each inner loop.

I have one more PR coming to optimize GPU NMS IR further.

please review @Laurawly @kevinthesun @zhiics @vinx13

[Laurawly]

LGTM

[masahi]

thanks @Laurawly

[mbrookhart]

Kudos!