masahi commented on a change in pull request #5116: Add thrust support for nms
URL: https://github.com/apache/incubator-tvm/pull/5116#discussion_r396164360
##########
File path: src/runtime/contrib/thrust/thrust.cu
##########
@@ -53,81 +54,112 @@ void thrust_sort(DLTensor* input,
thrust::copy(data_ptr, data_ptr + n_iter * n_values, values_ptr);
for (int i = 0 ; i < n_iter; ++i) {
- thrust::sequence(indices_ptr, indices_ptr + n_values);
+ int current_values = get_sort_len(i);
+ thrust::sequence(indices_ptr, indices_ptr + current_values);
if (is_ascend) {
- thrust::sort_by_key(values_ptr, values_ptr + n_values, indices_ptr);
+ thrust::sort_by_key(values_ptr, values_ptr + current_values,
indices_ptr);
} else {
- thrust::sort_by_key(values_ptr, values_ptr + n_values, indices_ptr,
+ thrust::sort_by_key(values_ptr, values_ptr + current_values, indices_ptr,
thrust::greater<DataType>());
}
- values_ptr += n_values;
- indices_ptr += n_values;
+ values_ptr += current_values;
+ indices_ptr += current_values;
}
}
-TVM_REGISTER_GLOBAL("tvm.contrib.thrust.sort")
-.set_body([](TVMArgs args, TVMRetValue* ret) {
- CHECK_GE(args.num_args, 4);
- DLTensor* input = args[0];
- DLTensor* values_out = args[1];
- DLTensor* indices_out = args[2];
- bool is_ascend = args[3];
-
- auto data_dtype = DLDataType2String(input->dtype);
- auto out_dtype = DLDataType2String(indices_out->dtype);
-
+template<typename F>
+void thrust_sort_common(DLTensor* input,
+ DLTensor* values_out,
+ DLTensor* indices_out,
+ bool is_ascend,
+ F &get_sort_len,
+ std::string data_dtype,
+ std::string out_dtype) {
if (data_dtype == "float32") {
if (out_dtype == "int32") {
- thrust_sort<float, int32_t>(input, values_out, indices_out, is_ascend);
+ thrust_sort<F, float, int32_t>(input, values_out, indices_out,
is_ascend, get_sort_len);
Review comment:
I think if you bring `typename F` after `typename, DataType, typename
IndicesType` in `thrust_sort`, you can remove explicit `F` in `thrust_sort<F,
...>`
----------------------------------------------------------------
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
