ptrendx commented on a change in pull request #16542: [WIP] Faster GPU NMS
optimization
URL: https://github.com/apache/incubator-mxnet/pull/16542#discussion_r338837623
##########
File path: src/operator/tensor/sort_op-inl.cuh
##########
@@ -246,15 +310,25 @@ inline typename
std::enable_if<(std::is_same<KDType,mshadow::half::half_t>::valu
SortByKeyImpl(mshadow::Tensor<gpu, 1, KDType> keys,
mshadow::Tensor<gpu, 1, VDType> values, bool is_ascend,
mshadow::Tensor<gpu, 1, char>* workspace,
- const int begin_bit, const int end_bit) {
+ const int begin_bit, const int end_bit,
+ mshadow::Tensor<gpu, 1, KDType>* sorted_keys,
+ mshadow::Tensor<gpu, 1, VDType>* sorted_values) {
CHECK_EQ(keys.CheckContiguous(), true);
CHECK_EQ(values.CheckContiguous(), true);
#if CUDA_VERSION >= 9000
cudaStream_t stream = mshadow::Stream<gpu>::GetStream(keys.stream_);
- thrust::device_ptr<__half> key_iter = thrust::device_pointer_cast(
- reinterpret_cast<__half*>(keys.dptr_));
- thrust::device_ptr<__half> value_iter = thrust::device_pointer_cast(
- reinterpret_cast<__half*>(values.dptr_));
+ auto* k = &keys;
+ auto* v = &values;
+ if (sorted_keys != nullptr) {
+ k = sorted_keys;
+ mshadow::Copy(*sorted_keys, keys, keys.stream_);
+ }
+ if (sorted_values != nullptr) {
+ v = sorted_values;
+ mshadow::Copy(*sorted_values, values, values.stream_);
+ }
+ const auto key_iter =
thrust::device_pointer_cast(reinterpret_cast<half*>(k->dptr_));
+ const auto value_iter =
thrust::device_pointer_cast(reinterpret_cast<half*>(v->dptr_));
Review comment:
Yup, a leftover from rebase.
----------------------------------------------------------------
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:
us...@infra.apache.org
With regards,
Apache Git Services
