Skip to content

Commit

Permalink
removed some WARs from /dispatch/*
Browse files Browse the repository at this point in the history 
added missing __host__ __device__ to tuple.inl
  • Loading branch information
@wnbell
wnbell committed Apr 26, 2010
1 parent cea016d commit 94ddf29
Show file tree
Hide file tree
Showing 4 changed files with 105 additions and 34 deletions.
77 changes: 77 additions & 0 deletions performance/min_index.test
View file
Original file line number Diff line number Diff line change
@@ -0,0 +1,77 @@
PREAMBLE = \
"""
#include <thrust/reduce.h>
#include <thrust/sequence.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>

using namespace thrust;

struct smaller_tuple
{
__host__ __device__
tuple<float,int> operator()(tuple<float,int> a, tuple<float,int> b)
{
if (a < b)
return a;
else
return b;
}
};

int min_index_slow(device_vector<float>& values)
{
device_vector<int> indices(values.size());
sequence(indices.begin(), indices.end());

tuple<float,int> init(values[0],0);

tuple<float,int> smallest = reduce(make_zip_iterator(make_tuple(values.begin(), indices.begin())),
make_zip_iterator(make_tuple(values.end(), indices.end())),
init,
smaller_tuple());
return get<1>(smallest);
}

int min_index_fast(device_vector<float>& values)
{
counting_iterator<int> begin(0);
counting_iterator<int> end(values.size());

tuple<float,int> init(values[0],0);

tuple<float,int> smallest = reduce(make_zip_iterator(make_tuple(values.begin(), begin)),
make_zip_iterator(make_tuple(values.end(), end)),
init,
smaller_tuple());
return get<1>(smallest);
}



"""

INITIALIZE = \
"""
thrust::host_vector<float> h_input = unittest::random_integers<float>($InputSize);
thrust::device_vector<float> d_input = h_input;

"""

TIME = \
"""
$Function(d_input);
"""

FINALIZE = \
"""
RECORD_TIME();
RECORD_THROUGHPUT(double($InputSize));
RECORD_BANDWIDTH(sizeof(float) * double($InputSize));
"""

Functions = ['min_index_slow','min_index_fast']
InputSizes = [2**22]

TestVariables = [('Function',Functions), ('InputSize', InputSizes)]

7 changes: 2 additions & 5 deletions thrust/detail/device/cuda/reduce_n.inl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -240,12 +240,9 @@ template<typename InputIterator,
return init;

// whether to perform blockwise reductions in shared memory or global memory
static const bool use_smem = sizeof(OutputType) <= 64;
thrust::detail::integral_constant<bool, sizeof(OutputType) <= 64> use_smem;

// XXX WAR nvcc 3.0 unused variable warning
(void) use_smem;

return detail::reduce_n(first, n, init, binary_op, thrust::detail::integral_constant<bool, use_smem>());
return detail::reduce_n(first, n, init, binary_op, use_smem);
} // end reduce_n()

} // end namespace cuda
Expand Down
38 changes: 9 additions & 29 deletions thrust/detail/device/dispatch/copy.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,13 +22,10 @@

namespace thrust
{

namespace detail
{

namespace device
{

namespace dispatch
{

Expand Down Expand Up @@ -71,35 +68,18 @@ template<typename InputIterator,
Space1,
Space2)
{
// // inspect both spaces
// XXX this is what we want
// typedef typename thrust::detail::integral_constant<bool,
// thrust::detail::is_convertible<Space1,thrust::detail::cuda_device_space_tag>::value ||
// thrust::detail::is_convertible<Space2,thrust::detail::cuda_device_space_tag>::value
// > is_one_of_the_spaces_cuda;

// XXX WAR nvcc 3.0b crash
static const bool temp =
thrust::detail::is_same<Space1,thrust::detail::cuda_device_space_tag>::value ||
thrust::detail::is_same<Space2,thrust::detail::cuda_device_space_tag>::value
;

typedef typename thrust::detail::eval_if<
temp,
thrust::detail::true_type,
thrust::detail::false_type
>::type is_one_of_the_spaces_cuda;
// inspect both spaces
typedef typename thrust::detail::integral_constant<bool,
thrust::detail::is_convertible<Space1,thrust::detail::cuda_device_space_tag>::value ||
thrust::detail::is_convertible<Space2,thrust::detail::cuda_device_space_tag>::value
> is_one_of_the_spaces_cuda;

return copy(first, last, result,
is_one_of_the_spaces_cuda());
} // end copy()


} // end dispatch

} // end device

} // end detail

} // end thrust
} // end namespace dispatch
} // end namespace device
} // end namespace detail
} // end namespace thrust

17 changes: 17 additions & 0 deletions thrust/detail/tuple.inl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -791,6 +791,7 @@ namespace detail
{

template<class T1, class T2>
__host__ __device__
inline bool eq(const T1& lhs, const T2& rhs) {
return lhs.get_head() == rhs.get_head() &&
eq(lhs.get_tail(), rhs.get_tail());
Expand All @@ -799,47 +800,57 @@ template<>
inline bool eq<null_type,null_type>(const null_type&, const null_type&) { return true; }

template<class T1, class T2>
__host__ __device__
inline bool neq(const T1& lhs, const T2& rhs) {
return lhs.get_head() != rhs.get_head() ||
neq(lhs.get_tail(), rhs.get_tail());
}
template<>
__host__ __device__
inline bool neq<null_type,null_type>(const null_type&, const null_type&) { return false; }

template<class T1, class T2>
__host__ __device__
inline bool lt(const T1& lhs, const T2& rhs) {
return lhs.get_head() < rhs.get_head() ||
!(rhs.get_head() < lhs.get_head()) &&
lt(lhs.get_tail(), rhs.get_tail());
}
template<>
__host__ __device__
inline bool lt<null_type,null_type>(const null_type&, const null_type&) { return false; }

template<class T1, class T2>
__host__ __device__
inline bool gt(const T1& lhs, const T2& rhs) {
return lhs.get_head() > rhs.get_head() ||
!(rhs.get_head() > lhs.get_head()) &&
gt(lhs.get_tail(), rhs.get_tail());
}
template<>
__host__ __device__
inline bool gt<null_type,null_type>(const null_type&, const null_type&) { return false; }

template<class T1, class T2>
__host__ __device__
inline bool lte(const T1& lhs, const T2& rhs) {
return lhs.get_head() <= rhs.get_head() &&
( !(rhs.get_head() <= lhs.get_head()) ||
lte(lhs.get_tail(), rhs.get_tail()));
}
template<>
__host__ __device__
inline bool lte<null_type,null_type>(const null_type&, const null_type&) { return true; }

template<class T1, class T2>
__host__ __device__
inline bool gte(const T1& lhs, const T2& rhs) {
return lhs.get_head() >= rhs.get_head() &&
( !(rhs.get_head() >= lhs.get_head()) ||
gte(lhs.get_tail(), rhs.get_tail()));
}
template<>
__host__ __device__
inline bool gte<null_type,null_type>(const null_type&, const null_type&) { return true; }

} // end detail
Expand All @@ -849,6 +860,7 @@ inline bool gte<null_type,null_type>(const null_type&, const null_type&) { retur
// equal ----

template<class T1, class T2, class S1, class S2>
__host__ __device__
inline bool operator==(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
{
// XXX support this eventually -jph
Expand All @@ -861,6 +873,7 @@ inline bool operator==(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S
// not equal -----

template<class T1, class T2, class S1, class S2>
__host__ __device__
inline bool operator!=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
{
// XXX support this eventually -jph
Expand All @@ -872,6 +885,7 @@ inline bool operator!=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S

// <
template<class T1, class T2, class S1, class S2>
__host__ __device__
inline bool operator<(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
{
// XXX support this eventually -jph
Expand All @@ -883,6 +897,7 @@ inline bool operator<(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2

// >
template<class T1, class T2, class S1, class S2>
__host__ __device__
inline bool operator>(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
{
// XXX support this eventually -jph
Expand All @@ -894,6 +909,7 @@ inline bool operator>(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2

// <=
template<class T1, class T2, class S1, class S2>
__host__ __device__
inline bool operator<=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
{
// XXX support this eventually -jph
Expand All @@ -905,6 +921,7 @@ inline bool operator<=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S

// >=
template<class T1, class T2, class S1, class S2>
__host__ __device__
inline bool operator>=(const detail::cons<T1, T2>& lhs, const detail::cons<S1, S2>& rhs)
{
// XXX support this eventually -jph
Expand Down

0 comments on commit 94ddf29

Please sign in to comment.