masahi commented on a change in pull request #5116: Add thrust support for nms
URL: https://github.com/apache/incubator-tvm/pull/5116#discussion_r396053175
##########
File path: src/runtime/contrib/thrust/thrust.cu
##########
@@ -65,6 +65,105 @@ void thrust_sort(DLTensor* input,
}
}
+// Performs sorting along axis -1 and returns both sorted values and indices.
+template<typename DataType, typename IndicesType>
+void thrust_sort_nms(DLTensor* input,
+ DLTensor* valid_count,
+ DLTensor* out_values,
+ DLTensor* out_indices,
+ bool is_ascend) {
+ thrust::device_ptr<IndicesType> valid_count_ptr(static_cast<IndicesType
*>(valid_count->data));
+ thrust::device_ptr<DataType> data_ptr(static_cast<DataType *>(input->data));
+ thrust::device_ptr<DataType> values_ptr(static_cast<DataType
*>(out_values->data));
+ thrust::device_ptr<IndicesType> indices_ptr(static_cast<IndicesType
*>(out_indices->data));
+
+ int n_values = input->shape[input->ndim - 1];
+ int n_iter = 1;
+ for (int i = 0; i < input->ndim - 1; ++i) {
+ n_iter *= input->shape[i];
+ }
+
+ thrust::copy(data_ptr, data_ptr + n_iter * n_values, values_ptr);
+
+ for (int i = 0 ; i < n_iter; ++i) {
+ int current_values = valid_count_ptr[i];
Review comment:
`thrust_sort_common` should take a function (lambda) as an argument, which
takes current iter (`i` in the code) and returns:
* For existing `thrust_sort`, always return `n_values`
* For `thrust_sort_nms`, returns `valid_count_ptr[i]` (`current_values`)
Does this make sense? I remember NVCC should not have any problem compiling
C++ lambdas.
----------------------------------------------------------------
This is an automated message from the Apache Git Service.
To respond to the message, please log on to GitHub and use the
URL above to go to the specific comment.
For queries about this service, please contact Infrastructure at:
[email protected]
With regards,
Apache Git Services
