masahi commented on a change in pull request #5116: Add thrust support for nms
URL: https://github.com/apache/incubator-tvm/pull/5116#discussion_r395939482
##########
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];
+ thrust::sequence(indices_ptr, indices_ptr + current_values);
+ if (is_ascend) {
+ thrust::sort_by_key(values_ptr, values_ptr + current_values,
indices_ptr);
+ } else {
+ thrust::sort_by_key(values_ptr, values_ptr + current_values, indices_ptr,
+ thrust::greater<DataType>());
+ }
+ values_ptr += current_values;
+ indices_ptr += current_values;
+ }
+}
+
+TVM_REGISTER_GLOBAL("tvm.contrib.thrust.sort_nms")
+.set_body([](TVMArgs args, TVMRetValue* ret) {
+ CHECK_GE(args.num_args, 5);
+ DLTensor* input = args[0];
+ DLTensor* valid_count = args[1];
+ DLTensor* values_out = args[2];
+ DLTensor* indices_out = args[3];
+ bool is_ascend = args[4];
+
+ auto data_dtype = DLDataType2String(input->dtype);
+ auto out_dtype = DLDataType2String(indices_out->dtype);
+
+ if (data_dtype == "float32") {
+ if (out_dtype == "int32") {
+ thrust_sort_nms<float, int32_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "int64") {
+ thrust_sort_nms<float, int64_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float32") {
+ thrust_sort_nms<float, float>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float64") {
+ thrust_sort_nms<float, double>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else if (data_dtype == "float64") {
+ if (out_dtype == "int32") {
+ thrust_sort_nms<double, int32_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "int64") {
+ thrust_sort_nms<double, int64_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float32") {
+ thrust_sort_nms<double, float>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float64") {
+ thrust_sort_nms<double, double>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else if (data_dtype == "int32") {
+ if (out_dtype == "int32") {
+ thrust_sort_nms<int32_t, int32_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "int64") {
+ thrust_sort_nms<int32_t, int64_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float32") {
+ thrust_sort_nms<int32_t, float>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float64") {
+ thrust_sort_nms<int32_t, double>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else if (data_dtype == "int64") {
+ if (out_dtype == "int32") {
+ thrust_sort_nms<int64_t, int32_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "int64") {
+ thrust_sort_nms<int64_t, int64_t>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float32") {
+ thrust_sort_nms<int64_t, float>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else if (out_dtype == "float64") {
+ thrust_sort_nms<int64_t, double>(input, valid_count, values_out,
indices_out, is_ascend);
+ } else {
+ LOG(FATAL) << "Unsupported output dtype: " << out_dtype;
+ }
+ } else {
+ LOG(FATAL) << "Unsupported input dtype: " << data_dtype;
+ }
+});
+
Review comment:
Likely we can share this if/else with `thrust_sort` using template or lambda
----------------------------------------------------------------
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
