tvm-commits mailing list archives

Site index · List index
Message view « Date » · « Thread »
Top « Date » · « Thread »
From GitBox <...@apache.org>
Subject [GitHub] [incubator-tvm] masahi commented on a change in pull request #5116: Add thrust support for nms
Date Sun, 22 Mar 2020 04:34:35 GMT
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:
users@infra.apache.org


With regards,
Apache Git Services

Mime
View raw message