Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

introduce get_iterator_value #783

Merged
jaredhoberock merged 8 commits into from
May 2, 2016
Merged

introduce get_iterator_value #783

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
3b79f66
get_iterator_value: safe pointer/iterator dereferencing
3gx Apr 29, 2016
718c996
bug fix
3gx Apr 29, 2016
2b8e2f2
get_iterator_value does not belong in thryst/system/detail/generic be…
3gx Apr 29, 2016
10c409d
add {min,max,minmax}_element testing with transform_iterator
3gx Apr 29, 2016
5b848c5
add {min,max,minmax}_element testing with raw device pointer and thru…
3gx Apr 29, 2016
7ee3fc9
remove duplicate #pragma once
3gx Apr 29, 2016
f9c9daa
typo fix
3gx Apr 29, 2016
c73d64a
invoke get_iterator_value via thrust::detail::get_iterator_value
3gx Apr 29, 2016
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 19 additions & 0 deletions testing/backend/cuda/max_element.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,3 +83,22 @@ void TestMaxElementCudaStreams()
}
DECLARE_UNITTEST(TestMaxElementCudaStreams);

void TestMaxElementDevicePointer()
{
typedef thrust::device_vector<int> Vector;
typedef typename Vector::value_type T;

Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;

T* raw_ptr = thrust::raw_pointer_cast(data.data());
size_t n = data.size();
ASSERT_EQUAL( thrust::max_element(thrust::device, raw_ptr, raw_ptr+n) - raw_ptr, 1);
ASSERT_EQUAL( thrust::max_element(thrust::device, raw_ptr, raw_ptr+n, thrust::greater<T>()) - raw_ptr, 2);
}
DECLARE_UNITTEST(TestMaxElementDevicePointer);
19 changes: 19 additions & 0 deletions testing/backend/cuda/min_element.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -83,3 +83,22 @@ void TestMinElementCudaStreams()
}
DECLARE_UNITTEST(TestMinElementCudaStreams);

void TestMinElementDevicePointer()
{
typedef thrust::device_vector<int> Vector;
typedef typename Vector::value_type T;

Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;

T* raw_ptr = thrust::raw_pointer_cast(data.data());
size_t n = data.size();
ASSERT_EQUAL( thrust::min_element(thrust::device, raw_ptr, raw_ptr+n) - raw_ptr, 2);
ASSERT_EQUAL( thrust::min_element(thrust::device, raw_ptr, raw_ptr+n, thrust::greater<T>()) - raw_ptr, 1);
}
DECLARE_UNITTEST(TestMinElementDevicePointer);
20 changes: 20 additions & 0 deletions testing/backend/cuda/minmax_element.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -102,3 +102,23 @@ void TestMinMaxElementCudaStreams()
}
DECLARE_UNITTEST(TestMinMaxElementCudaStreams);

void TestMinMaxElementDevicePointer()
{
typedef thrust::device_vector<int> Vector;
typedef typename Vector::value_type T;

Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;

T* raw_ptr = thrust::raw_pointer_cast(data.data());
size_t n = data.size();
ASSERT_EQUAL( thrust::minmax_element(thrust::device, raw_ptr, raw_ptr+n).first - raw_ptr, 2);
ASSERT_EQUAL( thrust::minmax_element(thrust::device, raw_ptr, raw_ptr+n).second - raw_ptr, 1);
}
DECLARE_UNITTEST(TestMinMaxElementDevicePointer);

26 changes: 26 additions & 0 deletions testing/max_element.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,8 @@
#include <unittest/unittest.h>
#include <thrust/extrema.h>
#include <thrust/iterator/retag.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/functional.h>

template <class Vector>
void TestMaxElementSimple(void)
Expand All @@ -23,6 +25,30 @@ void TestMaxElementSimple(void)
}
DECLARE_VECTOR_UNITTEST(TestMaxElementSimple);

template <class Vector>
void TestMaxElementWithTransform(void)
{
typedef typename Vector::value_type T;

Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;

ASSERT_EQUAL( *thrust::max_element(
thrust::make_transform_iterator(data.begin(), thrust::negate<T>()),
thrust::make_transform_iterator(data.end(), thrust::negate<T>())), -1);
ASSERT_EQUAL( *thrust::max_element(
thrust::make_transform_iterator(data.begin(), thrust::negate<T>()),
thrust::make_transform_iterator(data.end(), thrust::negate<T>()),
thrust::greater<T>()), -5);

}
DECLARE_VECTOR_UNITTEST(TestMaxElementWithTransform);

template<typename T>
void TestMaxElement(const size_t n)
{
Expand Down
24 changes: 24 additions & 0 deletions testing/min_element.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,30 @@ void TestMinElementSimple(void)
}
DECLARE_VECTOR_UNITTEST(TestMinElementSimple);

template <class Vector>
void TestMinElementWithTransform(void)
{
typedef typename Vector::value_type T;

Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;

ASSERT_EQUAL( *thrust::min_element(
thrust::make_transform_iterator(data.begin(), thrust::negate<T>()),
thrust::make_transform_iterator(data.end(), thrust::negate<T>())), -5);
ASSERT_EQUAL( *thrust::min_element(
thrust::make_transform_iterator(data.begin(), thrust::negate<T>()),
thrust::make_transform_iterator(data.end(), thrust::negate<T>()),
thrust::greater<T>()), -1);

}
DECLARE_VECTOR_UNITTEST(TestMinElementWithTransform);

template<typename T>
void TestMinElement(const size_t n)
{
Expand Down
23 changes: 23 additions & 0 deletions testing/minmax_element.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,29 @@ void TestMinMaxElementSimple(void)
ASSERT_EQUAL( thrust::minmax_element(data.begin(), data.end()).second - data.begin(), 1);
}
DECLARE_VECTOR_UNITTEST(TestMinMaxElementSimple);

template <class Vector>
void TestMinMaxElementWithTransform(void)
{
typedef typename Vector::value_type T;

Vector data(6);
data[0] = 3;
data[1] = 5;
data[2] = 1;
data[3] = 2;
data[4] = 5;
data[5] = 1;

ASSERT_EQUAL( *thrust::minmax_element(
thrust::make_transform_iterator(data.begin(), thrust::negate<T>()),
thrust::make_transform_iterator(data.end(), thrust::negate<T>())).first, -5);
ASSERT_EQUAL( *thrust::minmax_element(
thrust::make_transform_iterator(data.begin(), thrust::negate<T>()),
thrust::make_transform_iterator(data.end(), thrust::negate<T>())).second, -1);
}
DECLARE_VECTOR_UNITTEST(TestMinMaxElementWithTransform);


template<typename T>
void TestMinMaxElement(const size_t n)
Expand Down
49 changes: 49 additions & 0 deletions thrust/detail/get_iterator_value.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,49 @@
#pragma once
/*
* Copyright 2008-2016 NVIDIA Corporation
*
* Licensed under the Apache License, Version 2.0 (the "License");
* you may not use this file except in compliance with the License.
* You may obtain a copy of the License at
*
* http://www.apache.org/licenses/LICENSE-2.0
*
* Unless required by applicable law or agreed to in writing, software
* distributed under the License is distributed on an "AS IS" BASIS,
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
* See the License for the specific language governing permissions and
* limitations under the License.
*/

#include <thrust/detail/config.h>

namespace thrust {
namespace detail {

// get_iterator_value specialization on iterators
// --------------------------------------------------
// it is okay to dereference iterator in usual way
template<typename DerivedPolicy, typename Iterator>
__host__ __device__
typename thrust::iterator_traits<Iterator>::value_type
get_iterator_value(thrust::execution_policy<DerivedPolicy> &, Iterator it)
Copy link
Contributor

@jaredhoberock jaredhoberock May 2, 2016 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I'm not sure this implementation will work as intended because it ignores the execution policy. For example, if the user provides a CUDA stream in the execution policy, that will be ignored, and this dereference will use the default stream.

Wouldn't it be more correct for the body of get_iterator_value to do something like call thrust::copy?

Copy link
Contributor Author

@3gx 3gx May 2, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why wouldn't it, aren't true iterators supposed to be safe for usual derlerencing? If not, what is the interface to deference an iterator with an execution policy? This passes all the tests.

Copy link
Contributor

@jaredhoberock jaredhoberock May 2, 2016 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think that we intend for get_iterator_value() itself to be the interface to dereference an iterator with an execution policy. So, we have to build the implementation and ensure that it handles all cases correctly. I'm concerned that this implementation doesn't handle all cases correctly, namely, those cases where the user has provided a custom execution policy which is different from whatever is tagged in the iterator.

Copy link
Contributor Author

@3gx 3gx May 2, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not following. What would be an example of dereferencing an iterator with execution policy?

Copy link
Contributor Author

@3gx 3gx May 2, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

To add to this: it will be a non-issue because nothing nothing changes for iterators: they are still � dereference via 'operator*', as it was earlier. If users were hitting bugs with such dereferencing, we would have already known. We need a counter example to show it is unsafe now and worked before this change.

Copy link
Contributor

@jaredhoberock jaredhoberock May 2, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Suppose I call max_element(par.on(my_stream), vec.begin(), vec.end()) and look at the visual profiler timeline of my application. What sort of synchronization behavior would I observe with this implementation?

I believe I would see that all CUDA streams in my application synchronize at the call to max_element because the iterator dereference inside of get_iterator_value() will use the default stream instead of the stream contained inside of my execution policy.

Remember that in addition to preventing a crash due to dereferencing a raw device pointer on the host, we also have to ensure that the user's execution policy is used for this dereference. The one overload's use of get_value() ensures both of those things happen, but the other overload's plain iterator dereference does not.

I think that using thrust::copy will solve this problem.

Copy link
Contributor Author

@3gx 3gx May 2, 2016 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I understand. However, this fix doesn't change the previously existed behaviour, but fixes a bug when a raw pointer is passed to max_element.

That being said, replacing *it with

  typename thrust::iterator_traits<Iterator>::value_type value;
  thrust::copy(exec, it,it+1, &value); 
  return value;

makes the reproducer to die with

terminate called after throwing an instance of 'thrust::system::system_error'
  what():  an illegal memory access was encountered
Aborted (core dumped)

when used with transform_iterator and I wasn't able to quickly nail down the issue. I agree the get_value_iterator needs to be enhanced to respect customer user policy, but it may take more time.

I suggest either accepting this fix, or reverting #777 to make sure the user codes are not broken and then continue working on this PR to make get_iterator_value to respect user policy.

{
return *it;
} // get_iterator_value(exec,Iterator);

// get_iterator_value specialization on pointer
// ----------------------------------------------
// we can't just dereference a pointer in usual way, because
// it may point to a location in the device memory.
// we use get_value(exec,pointer*) function
// to perform a dereferencing consistent with the execution policy
template<typename DerivedPolicy, typename Pointer>
__host__ __device__
typename thrust::detail::pointer_traits<Pointer*>::element_type
get_iterator_value(thrust::execution_policy<DerivedPolicy> &exec, Pointer* ptr)
{
return get_value(derived_cast(exec),ptr);
} // get_iterator_value(exec,Pointer*)
Copy link
Contributor

@jaredhoberock jaredhoberock May 2, 2016 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It looks like this implementation only works for raw pointers. What will happen with types like thrust::device_ptr?

Copy link
Contributor Author

@3gx 3gx May 2, 2016

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Dereferencing device_ptr is safe on the host. Upon dereferencing on the host data is copied from the device, and so it is treated as an iterator. This specialisation is meant for raw pointers only that might be unsafe to dereference on the host.


} // namespace detail
} // namespace thrust
8 changes: 5 additions & 3 deletions thrust/system/detail/generic/extrema.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -22,6 +22,7 @@
#pragma once

#include <thrust/detail/config.h>
#include <thrust/detail/get_iterator_value.h>
#include <thrust/extrema.h>
#include <thrust/functional.h>
#include <thrust/pair.h>
Expand Down Expand Up @@ -172,7 +173,7 @@ ForwardIterator min_element(thrust::execution_policy<DerivedPolicy> &exec,
(exec,
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))),
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))) + (last - first),
thrust::tuple<InputType, IndexType>(get_value(derived_cast(exec), &first[0]), 0),
thrust::tuple<InputType, IndexType>(thrust::detail::get_iterator_value(derived_cast(exec), first), 0),
detail::min_element_reduction<InputType, IndexType, BinaryPredicate>(comp));

return first + thrust::get<1>(result);
Expand Down Expand Up @@ -209,7 +210,7 @@ ForwardIterator max_element(thrust::execution_policy<DerivedPolicy> &exec,
(exec,
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))),
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))) + (last - first),
thrust::tuple<InputType, IndexType>(get_value(derived_cast(exec),&first[0]), 0),
thrust::tuple<InputType, IndexType>(thrust::detail::get_iterator_value(derived_cast(exec),first), 0),
detail::max_element_reduction<InputType, IndexType, BinaryPredicate>(comp));

return first + thrust::get<1>(result);
Expand Down Expand Up @@ -247,7 +248,8 @@ thrust::pair<ForwardIterator,ForwardIterator> minmax_element(thrust::execution_p
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))),
thrust::make_zip_iterator(thrust::make_tuple(first, thrust::counting_iterator<IndexType>(0))) + (last - first),
detail::duplicate_tuple<InputType, IndexType>(),
detail::duplicate_tuple<InputType, IndexType>()(thrust::tuple<InputType, IndexType>(get_value(derived_cast(exec),&first[0]), 0)),
detail::duplicate_tuple<InputType, IndexType>()(
thrust::tuple<InputType, IndexType>(thrust::detail::get_iterator_value(derived_cast(exec),first), 0)),
detail::minmax_element_reduction<InputType, IndexType, BinaryPredicate>(comp));

return thrust::make_pair(first + thrust::get<1>(thrust::get<0>(result)), first + thrust::get<1>(thrust::get<1>(result)));
Expand Down