[BugFix] SSD fully supported on GPUs, updated deploy_ssd tutorial

[Laurawly]

Thanks to @vinx13 's pr #2420, argsort working now on GPUs.
Tested SSD full pipeline on NVIDIA K80c and Intel HD graphics. Performance improved compared with heterogenous results.
Please review @masahi @kevinthesun @zhiics

[masahi]

if I remember correctly, for opencl it should be tvm.opencl(0) or tvm.cl(0), isn't it?

[Laurawly]

Yes, sry I forgot to change.

[masahi]

@Laurawly @vinx13 can we share the sorting IR in this PR and @vinx13's PR #2420? They look identical.

[Laurawly]

I agree to put sort in a common file. And we can add a unitest for it as well.

[vinx13]

seems storage_sync is missing here, I will update my pr

[Laurawly]

@vinx13 Would you like to seperate argsort to a seperate file so that we can share the use of it? I can add unitest to it if needed.

[vinx13]

@Laurawly What's needed in ssd? Seems that you changed num_bbox in my pr to p_index[0], why only first element in p_index is used?

[kevinthesun]

Maybe we can make argsort a normal topi op? I'll add cpu implementation later.

[Laurawly]

@vinx13 p_index is the valid_count variable which is a 1D array resulted from the multibox operators. So instead of sorting all of data.shape[1] numbers, we only need to sort the first p_index[0] numbers.

[vinx13]

@Laurawly shouldn't be p_index[batch_id]? are you assuming batch = 1?

[Laurawly]

@vinx13 p_index only have one dimension. So it should be p_index[0].

[vinx13]

@kevinthesun @Laurawly The difficulty of sharing argsort (or extract it as a topi operator) is that we hope sort_num can be either a tvm.Tensor or constant array, but we can't use tvm.Expr to subscript a python array. Do you have ideas?

[vinx13]

@Laurawly still confused, if batch > 1, it should enter this if branch (since axis_mul_before * axis_mul_after > 1). Does p_index[tid] here mean that each batch has a different valid count?

[Laurawly]

@vinx13 From here https://github.com/dmlc/tvm/blob/master/topi/python/topi/cuda/nms.py#L368 axis is always 1, so axis_mul_before and axis_mul_after are both 1.

[vinx13]

@Laurawly since ndim == 2, axis == 1, the actual loop is like

for i in range(0, 2):
  if i < 1:
     axis_mul_before *= data.shape[i]

I assume that axis_mul_after == 1, axis_mul_before == data.shape[0], which is batch size, right?

[Laurawly]

@vinx13 Yeah, that's right. I see what you mean. So each batch could have a different valid count when batch_size > 1. I shouldn't have assumed batch_size = 1. I just pushed the changes.

[vinx13]

@Laurawly Btw have you checked the data race in nms ir? Seems __syncthreads and global barrier (maybe we should rewrite the ir to avoid global barrier) are needed on CUDA. I sometimes get incorrect nms results in my pr.

[Laurawly]

@vinx13 Does the conflict happen in argsort_ir?

[vinx13]

@Laurawly the conflict happens in nms_ir, I replaced blockIdx with vthread and added storage_sync and it worked, but my current solution is not efficient

[Laurawly]

@vinx13 I don't see conflicts in my nms_ir using blockIdx.x, but I'll double check. Why do you want to replace blockIdx.x with vthread?

[vinx13]

@Laurawly If the data written by other threads is needed (probably this line if_scope(p_out[b * num_anchors * 6 + offset_l] >= 0), there may be data race because lack of synchronization.

[Laurawly]

@vinx13 There's no data conflict for p_out in SSD because nms_topk = -1. For the line you mentioned, because writing and reading of p_out happen in the same thread block (if_scope(p_out[b * num_anchors * 6 + offset_i] >= 0) and p_out[b * num_anchors * 6 + offset_i] = -1.0)

[vinx13]

@Laurawly the writing p_out[base_idx + offset_i] = -1.0 is in thread i, while the reading of p_out[base_idx + offset_l] in ib.if_scope(p_out[base_idx + offset_l] >= 0) and p_out[base_idx + offset_l] == p_out[b * num_anchors * 6 + offset_i])) is in all threads. And there is no synchrozation after each iteration of for_range(0, p_valid_count[b]) as l, is there confilict in this case?

[Laurawly]

@vinx13 No, because there's a condition that i > l. And because you iterate via l sequentially, when you read from p_out[base_idx + offset_l], the writing of it should be already finished. Such as when l==0, for thread 0, it couldn't write to p_out[base_idx + offset_0] so it returned; for thread 1 to thread n, they need to read from p_out[base_idx + offset_0] but we have already know thread 0 won't write to it. When l==1, for thread 0, 1, they won't write to p_out, for thread 2, it may need to read p_out[base_idx + offset_1] before thread 1 finish writing to it in iteration l==0. But in that case, it means p_out[base_idx + offset_1] == -1 and p_out[base_idx + offset_2] != -1 because thread 2 finishes much earlier than 'thread 1'. So p_out[base_idx + offset_l] ==p_out[base_idx + offset_i] won't be true either way.

[vinx13]

@Laurawly I see, thanks for your clarification

[masahi]

thanks @Laurawly @vinx13 @kevinthesun @zhiics this is merged.