tvm-commits mailing list archives

Site index · List index
Message view « Date » · « Thread »
Top « Date » · « Thread »
From GitBox <...@apache.org>
Subject [GitHub] [incubator-tvm] Laurawly commented on a change in pull request #5116: Add thrust support for nms
Date Sun, 22 Mar 2020 04:07:16 GMT
Laurawly commented on a change in pull request #5116: Add thrust support for nms
URL: https://github.com/apache/incubator-tvm/pull/5116#discussion_r396052230
 
 

 ##########
 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:
   Could you elaborate more on this?

----------------------------------------------------------------
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