ICODE-ADC/lammps
4
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update Kokkos library to r2.6.00

Browse Source
This commit is contained in:
Stan Moore
2018-03-08 10:57:08 -07:00
parent 0c4c002f34
commit 39786b1740
694 changed files with 12261 additions and 6745 deletions
Download Patch File Download Diff File Expand all files Collapse all files

150
lib/kokkos/algorithms/src/Kokkos_Random.hpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -1530,7 +1530,7 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,1,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0()))
if(idx<static_cast<IndexType>(a.extent(0)))
a(idx) = Rand::draw(gen,range);
}
rand_pool.free_state(gen);
@ -1555,8 +1555,8 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,2,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
a(idx,k) = Rand::draw(gen,range);
}
}
@ -1583,9 +1583,9 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,3,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
a(idx,k,l) = Rand::draw(gen,range);
}
}
@ -1611,10 +1611,10 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,4, IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
a(idx,k,l,m) = Rand::draw(gen,range);
}
}
@ -1640,11 +1640,11 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,5,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
a(idx,k,l,m,n) = Rand::draw(gen,range);
}
}
@ -1670,12 +1670,12 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,6,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_5());o++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(5));o++)
a(idx,k,l,m,n,o) = Rand::draw(gen,range);
}
}
@ -1701,13 +1701,13 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,7,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_5());o++)
for(IndexType p=0;p<static_cast<IndexType>(a.dimension_6());p++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(5));o++)
for(IndexType p=0;p<static_cast<IndexType>(a.extent(6));p++)
a(idx,k,l,m,n,o,p) = Rand::draw(gen,range);
}
}
@ -1733,14 +1733,14 @@ struct fill_random_functor_range<ViewType,RandomPool,loops,8,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_5());o++)
for(IndexType p=0;p<static_cast<IndexType>(a.dimension_6());p++)
for(IndexType q=0;q<static_cast<IndexType>(a.dimension_7());q++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(5));o++)
for(IndexType p=0;p<static_cast<IndexType>(a.extent(6));p++)
for(IndexType q=0;q<static_cast<IndexType>(a.extent(7));q++)
a(idx,k,l,m,n,o,p,q) = Rand::draw(gen,range);
}
}
@ -1765,7 +1765,7 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,1,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0()))
if(idx<static_cast<IndexType>(a.extent(0)))
a(idx) = Rand::draw(gen,begin,end);
}
rand_pool.free_state(gen);
@ -1790,8 +1790,8 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,2,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
a(idx,k) = Rand::draw(gen,begin,end);
}
}
@ -1818,9 +1818,9 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,3,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
a(idx,k,l) = Rand::draw(gen,begin,end);
}
}
@ -1846,10 +1846,10 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,4,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
a(idx,k,l,m) = Rand::draw(gen,begin,end);
}
}
@ -1875,11 +1875,11 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,5,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())){
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_1());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_2());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_3());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_4());o++)
if(idx<static_cast<IndexType>(a.extent(0))){
for(IndexType l=0;l<static_cast<IndexType>(a.extent(1));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(2));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(3));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(4));o++)
a(idx,l,m,n,o) = Rand::draw(gen,begin,end);
}
}
@ -1905,12 +1905,12 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,6,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_5());o++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(5));o++)
a(idx,k,l,m,n,o) = Rand::draw(gen,begin,end);
}
}
@ -1937,13 +1937,13 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,7,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_5());o++)
for(IndexType p=0;p<static_cast<IndexType>(a.dimension_6());p++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(5));o++)
for(IndexType p=0;p<static_cast<IndexType>(a.extent(6));p++)
a(idx,k,l,m,n,o,p) = Rand::draw(gen,begin,end);
}
}
@ -1969,14 +1969,14 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,8,IndexType>{
typename RandomPool::generator_type gen = rand_pool.get_state();
for(IndexType j=0;j<loops;j++) {
const IndexType idx = i*loops+j;
if(idx<static_cast<IndexType>(a.dimension_0())) {
for(IndexType k=0;k<static_cast<IndexType>(a.dimension_1());k++)
for(IndexType l=0;l<static_cast<IndexType>(a.dimension_2());l++)
for(IndexType m=0;m<static_cast<IndexType>(a.dimension_3());m++)
for(IndexType n=0;n<static_cast<IndexType>(a.dimension_4());n++)
for(IndexType o=0;o<static_cast<IndexType>(a.dimension_5());o++)
for(IndexType p=0;p<static_cast<IndexType>(a.dimension_6());p++)
for(IndexType q=0;q<static_cast<IndexType>(a.dimension_7());q++)
if(idx<static_cast<IndexType>(a.extent(0))) {
for(IndexType k=0;k<static_cast<IndexType>(a.extent(1));k++)
for(IndexType l=0;l<static_cast<IndexType>(a.extent(2));l++)
for(IndexType m=0;m<static_cast<IndexType>(a.extent(3));m++)
for(IndexType n=0;n<static_cast<IndexType>(a.extent(4));n++)
for(IndexType o=0;o<static_cast<IndexType>(a.extent(5));o++)
for(IndexType p=0;p<static_cast<IndexType>(a.extent(6));p++)
for(IndexType q=0;q<static_cast<IndexType>(a.extent(7));q++)
a(idx,k,l,m,n,o,p,q) = Rand::draw(gen,begin,end);
}
}
@ -1988,14 +1988,14 @@ struct fill_random_functor_begin_end<ViewType,RandomPool,loops,8,IndexType>{
template<class ViewType, class RandomPool, class IndexType = int64_t>
void fill_random(ViewType a, RandomPool g, typename ViewType::const_value_type range) {
int64_t LDA = a.dimension_0();
int64_t LDA = a.extent(0);
if(LDA>0)
parallel_for((LDA+127)/128,Impl::fill_random_functor_range<ViewType,RandomPool,128,ViewType::Rank,IndexType>(a,g,range));
}
template<class ViewType, class RandomPool, class IndexType = int64_t>
void fill_random(ViewType a, RandomPool g, typename ViewType::const_value_type begin,typename ViewType::const_value_type end ) {
int64_t LDA = a.dimension_0();
int64_t LDA = a.extent(0);
if(LDA>0)
parallel_for((LDA+127)/128,Impl::fill_random_functor_begin_end<ViewType,RandomPool,128,ViewType::Rank,IndexType>(a,g,begin,end));
}

48
lib/kokkos/algorithms/src/Kokkos_Sort.hpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -120,7 +120,6 @@ public:
KOKKOS_INLINE_FUNCTION
void operator() (const int& i) const {
// printf("copy: dst(%i) src(%i)\n",i+dst_offset,i);
copy_op::copy(dst_values,i+dst_offset,src_values,i);
}
};
@ -151,20 +150,22 @@ public:
DstViewType dst_values ;
perm_view_type sort_order ;
src_view_type src_values ;
int src_offset ;
copy_permute_functor( DstViewType const & dst_values_
, PermuteViewType const & sort_order_
, SrcViewType const & src_values_
, int const & src_offset_
)
: dst_values( dst_values_ )
, sort_order( sort_order_ )
, src_values( src_values_ )
, src_offset( src_offset_ )
{}
KOKKOS_INLINE_FUNCTION
void operator() (const int& i) const {
// printf("copy_permute: dst(%i) src(%i)\n",i,sort_order(i));
copy_op::copy(dst_values,i,src_values,sort_order(i));
copy_op::copy(dst_values,i,src_values,src_offset+sort_order(i));
}
};
@ -259,19 +260,21 @@ public:
// Create the permutation vector, the bin_offset array and the bin_count array. Can be called again if keys changed
void create_permute_vector() {
const size_t len = range_end - range_begin ;
Kokkos::parallel_for (Kokkos::RangePolicy<execution_space,bin_count_tag> (0,len),*this);
Kokkos::parallel_scan(Kokkos::RangePolicy<execution_space,bin_offset_tag> (0,bin_op.max_bins()) ,*this);
Kokkos::parallel_for ("Kokkos::Sort::BinCount",Kokkos::RangePolicy<execution_space,bin_count_tag> (0,len),*this);
Kokkos::parallel_scan("Kokkos::Sort::BinOffset",Kokkos::RangePolicy<execution_space,bin_offset_tag> (0,bin_op.max_bins()) ,*this);
Kokkos::deep_copy(bin_count_atomic,0);
Kokkos::parallel_for (Kokkos::RangePolicy<execution_space,bin_binning_tag> (0,len),*this);
Kokkos::parallel_for ("Kokkos::Sort::BinBinning",Kokkos::RangePolicy<execution_space,bin_binning_tag> (0,len),*this);
if(sort_within_bins)
Kokkos::parallel_for (Kokkos::RangePolicy<execution_space,bin_sort_bins_tag>(0,bin_op.max_bins()) ,*this);
Kokkos::parallel_for ("Kokkos::Sort::BinSort",Kokkos::RangePolicy<execution_space,bin_sort_bins_tag>(0,bin_op.max_bins()) ,*this);
}
// Sort a view with respect ot the first dimension using the permutation array
// Sort a subset of a view with respect to the first dimension using the permutation array
template<class ValuesViewType>
void sort( ValuesViewType const & values)
void sort( ValuesViewType const & values
, int values_range_begin
, int values_range_end) const
{
typedef
Kokkos::View< typename ValuesViewType::data_type,
@ -280,6 +283,10 @@ public:
scratch_view_type ;
const size_t len = range_end - range_begin ;
const size_t values_len = values_range_end - values_range_begin ;
if (len != values_len) {
Kokkos::abort("BinSort::sort: values range length != permutation vector length");
}
scratch_view_type
sorted_values("Scratch",
@ -297,19 +304,25 @@ public:
, offset_type /* PermuteViewType */
, ValuesViewType /* SrcViewType */
>
functor( sorted_values , sort_order , values );
functor( sorted_values , sort_order , values, values_range_begin - range_begin );
parallel_for( Kokkos::RangePolicy<execution_space>(0,len),functor);
parallel_for("Kokkos::Sort::CopyPermute", Kokkos::RangePolicy<execution_space>(0,len),functor);
}
{
copy_functor< ValuesViewType , scratch_view_type >
functor( values , range_begin , sorted_values );
parallel_for( Kokkos::RangePolicy<execution_space>(0,len),functor);
parallel_for("Kokkos::Sort::Copy", Kokkos::RangePolicy<execution_space>(0,len),functor);
}
}
template<class ValuesViewType>
void sort( ValuesViewType const & values ) const
{
this->sort( values, 0, /*values.extent(0)*/ range_end - range_begin );
}
// Get the permutation vector
KOKKOS_INLINE_FUNCTION
offset_type get_permute_vector() const { return sort_order;}
@ -327,7 +340,7 @@ public:
KOKKOS_INLINE_FUNCTION
void operator() (const bin_count_tag& tag, const int& i) const {
const int j = range_begin + i ;
bin_count_atomic(bin_op.bin(keys,j))++;
bin_count_atomic(bin_op.bin(keys, j))++;
}
KOKKOS_INLINE_FUNCTION
@ -512,7 +525,7 @@ void sort( ViewType const & view , bool const always_use_kokkos_sort = false)
Kokkos::Experimental::MinMaxScalar<typename ViewType::non_const_value_type> result;
Kokkos::Experimental::MinMax<typename ViewType::non_const_value_type> reducer(result);
parallel_reduce(Kokkos::RangePolicy<typename ViewType::execution_space>(0,view.extent(0)),
parallel_reduce("Kokkos::Sort::FindExtent",Kokkos::RangePolicy<typename ViewType::execution_space>(0,view.extent(0)),
Impl::min_max_functor<ViewType>(view),reducer);
if(result.min_val == result.max_val) return;
BinSort<ViewType, CompType> bin_sort(view,CompType(view.extent(0)/2,result.min_val,result.max_val),true);
@ -532,7 +545,7 @@ void sort( ViewType view
Kokkos::Experimental::MinMaxScalar<typename ViewType::non_const_value_type> result;
Kokkos::Experimental::MinMax<typename ViewType::non_const_value_type> reducer(result);
parallel_reduce( range_policy( begin , end )
parallel_reduce("Kokkos::Sort::FindExtent", range_policy( begin , end )
, Impl::min_max_functor<ViewType>(view),reducer );
if(result.min_val == result.max_val) return;
@ -541,8 +554,9 @@ void sort( ViewType view
bin_sort(view,begin,end,CompType((end-begin)/2,result.min_val,result.max_val),true);
bin_sort.create_permute_vector();
bin_sort.sort(view);
bin_sort.sort(view,begin,end);
}
}
#endif

7
lib/kokkos/algorithms/unit_tests/TestCuda.cpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -61,14 +61,9 @@ class cuda : public ::testing::Test {
protected:
static void SetUpTestCase()
{
std::cout << std::setprecision(5) << std::scientific;
Kokkos::HostSpace::execution_space::initialize();
Kokkos::Cuda::initialize( Kokkos::Cuda::SelectDevice(0) );
}
static void TearDownTestCase()
{
Kokkos::Cuda::finalize();
Kokkos::HostSpace::execution_space::finalize();
}
};

17
lib/kokkos/algorithms/unit_tests/TestOpenMP.cpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -60,25 +60,10 @@ protected:
static void SetUpTestCase()
{
std::cout << std::setprecision(5) << std::scientific;
int threads_count = 0;
#pragma omp parallel
{
#pragma omp atomic
++threads_count;
}
if (threads_count > 3) {
threads_count /= 2;
}
Kokkos::OpenMP::initialize( threads_count );
Kokkos::OpenMP::print_configuration( std::cout );
}
static void TearDownTestCase()
{
Kokkos::OpenMP::finalize();
}
};

6
lib/kokkos/algorithms/unit_tests/TestROCm.cpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -62,13 +62,9 @@ protected:
static void SetUpTestCase()
{
std::cout << std::setprecision(5) << std::scientific;
Kokkos::HostSpace::execution_space::initialize();
Kokkos::Experimental::ROCm::initialize( Kokkos::Experimental::ROCm::SelectDevice(0) );
}
static void TearDownTestCase()
{
Kokkos::Experimental::ROCm::finalize();
Kokkos::HostSpace::execution_space::finalize();
}
};

2
lib/kokkos/algorithms/unit_tests/TestRandom.hpp
View File

@ -34,7 +34,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER

5
lib/kokkos/algorithms/unit_tests/TestSerial.cpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -62,13 +62,10 @@ class serial : public ::testing::Test {
protected:
static void SetUpTestCase()
{
std::cout << std::setprecision (5) << std::scientific;
Kokkos::Serial::initialize ();
}
static void TearDownTestCase ()
{
Kokkos::Serial::finalize ();
}
};

104
lib/kokkos/algorithms/unit_tests/TestSort.hpp
View File

@ -34,7 +34,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -171,10 +171,10 @@ void test_3D_sort(unsigned int n) {
double sum_after = 0.0;
unsigned int sort_fails = 0;
Kokkos::parallel_reduce(keys.dimension_0(),sum3D<ExecutionSpace, KeyType>(keys),sum_before);
Kokkos::parallel_reduce(keys.extent(0),sum3D<ExecutionSpace, KeyType>(keys),sum_before);
int bin_1d = 1;
while( bin_1d*bin_1d*bin_1d*4< (int) keys.dimension_0() ) bin_1d*=2;
while( bin_1d*bin_1d*bin_1d*4< (int) keys.extent(0) ) bin_1d*=2;
int bin_max[3] = {bin_1d,bin_1d,bin_1d};
typename KeyViewType::value_type min[3] = {0,0,0};
typename KeyViewType::value_type max[3] = {100,100,100};
@ -186,8 +186,8 @@ void test_3D_sort(unsigned int n) {
Sorter.create_permute_vector();
Sorter.template sort< KeyViewType >(keys);
Kokkos::parallel_reduce(keys.dimension_0(),sum3D<ExecutionSpace, KeyType>(keys),sum_after);
Kokkos::parallel_reduce(keys.dimension_0()-1,bin3d_is_sorted_struct<ExecutionSpace, KeyType>(keys,bin_1d,min[0],max[0]),sort_fails);
Kokkos::parallel_reduce(keys.extent(0),sum3D<ExecutionSpace, KeyType>(keys),sum_after);
Kokkos::parallel_reduce(keys.extent(0)-1,bin3d_is_sorted_struct<ExecutionSpace, KeyType>(keys,bin_1d,min[0],max[0]),sort_fails);
double ratio = sum_before/sum_after;
double epsilon = 1e-10;
@ -205,24 +205,13 @@ void test_3D_sort(unsigned int n) {
template<class ExecutionSpace, typename KeyType>
void test_dynamic_view_sort(unsigned int n )
{
typedef typename ExecutionSpace::memory_space memory_space ;
typedef Kokkos::Experimental::DynamicView<KeyType*,ExecutionSpace> KeyDynamicViewType;
typedef Kokkos::View<KeyType*,ExecutionSpace> KeyViewType;
const size_t upper_bound = 2 * n ;
const size_t min_chunk_size = 1024;
const size_t total_alloc_size = n * sizeof(KeyType) * 1.2 ;
const size_t superblock_size = std::min(total_alloc_size, size_t(1000000));
typename KeyDynamicViewType::memory_pool
pool( memory_space()
, n * sizeof(KeyType) * 1.2
, 500 /* min block size in bytes */
, 30000 /* max block size in bytes */
, superblock_size
);
KeyDynamicViewType keys("Keys",pool,upper_bound);
KeyDynamicViewType keys("Keys", min_chunk_size, upper_bound);
keys.resize_serial(n);
@ -230,13 +219,15 @@ void test_dynamic_view_sort(unsigned int n )
// Test sorting array with all numbers equal
Kokkos::deep_copy(keys_view,KeyType(1));
Kokkos::Experimental::deep_copy(keys,keys_view);
Kokkos::deep_copy(keys,keys_view);
Kokkos::sort(keys, 0 /* begin */ , n /* end */ );
Kokkos::Random_XorShift64_Pool<ExecutionSpace> g(1931);
Kokkos::fill_random(keys_view,g,Kokkos::Random_XorShift64_Pool<ExecutionSpace>::generator_type::MAX_URAND);
Kokkos::Experimental::deep_copy(keys,keys_view);
ExecutionSpace::fence();
Kokkos::deep_copy(keys,keys_view);
//ExecutionSpace::fence();
double sum_before = 0.0;
double sum_after = 0.0;
@ -246,7 +237,9 @@ void test_dynamic_view_sort(unsigned int n )
Kokkos::sort(keys, 0 /* begin */ , n /* end */ );
Kokkos::Experimental::deep_copy( keys_view , keys );
ExecutionSpace::fence(); // Need this fence to prevent BusError with Cuda
Kokkos::deep_copy( keys_view , keys );
//ExecutionSpace::fence();
Kokkos::parallel_reduce(n,sum<ExecutionSpace, KeyType>(keys_view),sum_after);
Kokkos::parallel_reduce(n-1,is_sorted_struct<ExecutionSpace, KeyType>(keys_view),sort_fails);
@ -269,6 +262,74 @@ void test_dynamic_view_sort(unsigned int n )
//----------------------------------------------------------------------------
template<class ExecutionSpace>
void test_issue_1160()
{
Kokkos::View<int*, ExecutionSpace> element_("element", 10);
Kokkos::View<double*, ExecutionSpace> x_("x", 10);
Kokkos::View<double*, ExecutionSpace> v_("y", 10);
auto h_element = Kokkos::create_mirror_view(element_);
auto h_x = Kokkos::create_mirror_view(x_);
auto h_v = Kokkos::create_mirror_view(v_);
h_element(0) = 9;
h_element(1) = 8;
h_element(2) = 7;
h_element(3) = 6;
h_element(4) = 5;
h_element(5) = 4;
h_element(6) = 3;
h_element(7) = 2;
h_element(8) = 1;
h_element(9) = 0;
for (int i = 0; i < 10; ++i) {
h_v.access(i, 0) = h_x.access(i, 0) = double(h_element(i));
}
Kokkos::deep_copy(element_, h_element);
Kokkos::deep_copy(x_, h_x);
Kokkos::deep_copy(v_, h_v);
typedef decltype(element_) KeyViewType;
typedef Kokkos::BinOp1D< KeyViewType > BinOp;
int begin = 3;
int end = 8;
auto max = h_element(begin);
auto min = h_element(end - 1);
BinOp binner(end - begin, min, max);
Kokkos::BinSort<KeyViewType , BinOp > Sorter(element_,begin,end,binner,false);
Sorter.create_permute_vector();
Sorter.sort(element_,begin,end);
Sorter.sort(x_,begin,end);
Sorter.sort(v_,begin,end);
Kokkos::deep_copy(h_element, element_);
Kokkos::deep_copy(h_x, x_);
Kokkos::deep_copy(h_v, v_);
ASSERT_EQ(h_element(0), 9);
ASSERT_EQ(h_element(1), 8);
ASSERT_EQ(h_element(2), 7);
ASSERT_EQ(h_element(3), 2);
ASSERT_EQ(h_element(4), 3);
ASSERT_EQ(h_element(5), 4);
ASSERT_EQ(h_element(6), 5);
ASSERT_EQ(h_element(7), 6);
ASSERT_EQ(h_element(8), 1);
ASSERT_EQ(h_element(9), 0);
for (int i = 0; i < 10; ++i) {
ASSERT_EQ(h_element(i), int(h_x.access(i, 0)));
ASSERT_EQ(h_element(i), int(h_v.access(i, 0)));
}
}
//----------------------------------------------------------------------------
template<class ExecutionSpace, typename KeyType>
void test_sort(unsigned int N)
{
@ -278,6 +339,7 @@ void test_sort(unsigned int N)
test_3D_sort<ExecutionSpace,KeyType>(N);
test_dynamic_view_sort<ExecutionSpace,KeyType>(N*N);
#endif
test_issue_1160<ExecutionSpace>();
}
}

17
lib/kokkos/algorithms/unit_tests/TestThreads.cpp
View File

@ -35,7 +35,7 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
@ -63,25 +63,10 @@ protected:
static void SetUpTestCase()
{
std::cout << std::setprecision(5) << std::scientific;
unsigned num_threads = 4;
if (Kokkos::hwloc::available()) {
num_threads = Kokkos::hwloc::get_available_numa_count()
* Kokkos::hwloc::get_available_cores_per_numa()
// * Kokkos::hwloc::get_available_threads_per_core()
;
}
std::cout << "Threads: " << num_threads << std::endl;
Kokkos::Threads::initialize( num_threads );
}
static void TearDownTestCase()
{
Kokkos::Threads::finalize();
}
};

8
lib/kokkos/algorithms/unit_tests/UnitTestMain.cpp
View File

@ -35,16 +35,20 @@
// NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
// SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
//
// Questions? Contact H. Carter Edwards (hcedwar@sandia.gov)
// Questions? Contact Christian R. Trott (crtrott@sandia.gov)
//
// ************************************************************************
//@HEADER
*/
#include <gtest/gtest.h>
#include <Kokkos_Core.hpp>
int main(int argc, char *argv[]) {
Kokkos::initialize(argc,argv);
::testing::InitGoogleTest(&argc,argv);
return RUN_ALL_TESTS();
int result = RUN_ALL_TESTS();
Kokkos::finalize();
return result;
}