/*
 *  Copyright 2008-2013 NVIDIA Corporation
 *  Modifications Copyright© 2019 Advanced Micro Devices, Inc. All rights reserved.
 *
 *  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/functional.h>
#include <thrust/sort.h>

#include "test_header.hpp"

template <typename T>
struct less_div_10
{
    __host__ __device__ bool operator()(const T& lhs, const T& rhs) const
    {
        return ((int)lhs) / 10 < ((int)rhs) / 10;
    }
};

template <typename T>
struct greater_div_10
{
    __host__ __device__ bool operator()(const T& lhs, const T& rhs) const
    {
        return ((int)lhs) / 10 > ((int)rhs) / 10;
    }
};

template <typename T, unsigned int N>
void _TestStableSortByKeyWithLargeKeys(void)
{
    size_t n = (128 * 1024) / sizeof(FixedVector<T, N>);

    thrust::host_vector<FixedVector<T, N>> h_keys(n);
    thrust::host_vector<unsigned int>      h_vals(n);

    for(size_t i = 0; i < n; i++)
    {
        h_keys[i] = FixedVector<T, N>(rand());
        h_vals[i] = i;
    }

    thrust::device_vector<FixedVector<T, N>> d_keys = h_keys;
    thrust::device_vector<unsigned int>      d_vals = h_vals;

    thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());

    ASSERT_EQ_QUIET(h_keys, d_keys);
    ASSERT_EQ_QUIET(h_vals, d_vals);
}

TEST(StableSortByKeyLargeTests, TestStableSortByKeyWithLargeKeys)
{
    SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
    
    _TestStableSortByKeyWithLargeKeys<int, 4>();
    _TestStableSortByKeyWithLargeKeys<int, 8>();
    // XXX these take too long to compile
    _TestStableSortByKeyWithLargeKeys<int, 16>();
    _TestStableSortByKeyWithLargeKeys<int, 32>();
    _TestStableSortByKeyWithLargeKeys<int, 64>();
    //    _TestStableSortByKeyWithLargeKeys<int,  128>();
    //    _TestStableSortByKeyWithLargeKeys<int,  256>();
    //    _TestStableSortByKeyWithLargeKeys<int,  512>();
    //    _TestStableSortByKeyWithLargeKeys<int, 1024>();
    //    _TestStableSortByKeyWithLargeKeys<int, 2048>();
    //    _TestStableSortByKeyWithLargeKeys<int, 4096>();
    //    _TestStableSortByKeyWithLargeKeys<int, 8192>();
}

template <typename T, unsigned int N>
void _TestStableSortByKeyWithLargeValues(void)
{
    size_t n = (128 * 1024) / sizeof(FixedVector<T, N>);

    thrust::host_vector<unsigned int>      h_keys(n);
    thrust::host_vector<FixedVector<T, N>> h_vals(n);

    for(size_t i = 0; i < n; i++)
    {
        h_keys[i] = rand();
        h_vals[i] = FixedVector<T, N>(i);
    }

    thrust::device_vector<unsigned int>      d_keys = h_keys;
    thrust::device_vector<FixedVector<T, N>> d_vals = h_vals;

    thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());

    ASSERT_EQ_QUIET(h_keys, d_keys);
    ASSERT_EQ_QUIET(h_vals, d_vals);

    // so cuda::stable_merge_sort_by_key() is called
    thrust::stable_sort_by_key(
        h_keys.begin(), h_keys.end(), h_vals.begin(), greater_div_10<unsigned int>());
    thrust::stable_sort_by_key(
        d_keys.begin(), d_keys.end(), d_vals.begin(), greater_div_10<unsigned int>());

    ASSERT_EQ_QUIET(h_keys, d_keys);
    ASSERT_EQ_QUIET(h_vals, d_vals);
}

TEST(StableSortByKeyLargeTests, TestStableSortByKeyWithLargeValues)
{
    SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
    
    _TestStableSortByKeyWithLargeValues<int, 4>();
    _TestStableSortByKeyWithLargeValues<int, 8>();
    // XXX these take too long to compile
    _TestStableSortByKeyWithLargeValues<int, 16>();
    _TestStableSortByKeyWithLargeValues<int, 32>();
    _TestStableSortByKeyWithLargeValues<int, 64>();
    //    _TestStableSortByKeyWithLargeValues<int,  128>();
    //    _TestStableSortByKeyWithLargeValues<int,  256>();
    //    _TestStableSortByKeyWithLargeValues<int,  512>();
    //    _TestStableSortByKeyWithLargeValues<int, 1024>();
    //    _TestStableSortByKeyWithLargeValues<int, 2048>();
    //    _TestStableSortByKeyWithLargeValues<int, 4096>();
    //    _TestStableSortByKeyWithLargeValues<int, 8192>();
}

template <typename T, unsigned int N>
void _TestStableSortByKeyWithLargeKeysAndValues(void)
{
    size_t n = (128 * 1024) / sizeof(FixedVector<T, N>);

    thrust::host_vector<FixedVector<T, N>> h_keys(n);
    thrust::host_vector<FixedVector<T, N>> h_vals(n);

    for(size_t i = 0; i < n; i++)
    {
        h_keys[i] = FixedVector<T, N>(rand());
        h_vals[i] = FixedVector<T, N>(i);
    }

    thrust::device_vector<FixedVector<T, N>> d_keys = h_keys;
    thrust::device_vector<FixedVector<T, N>> d_vals = h_vals;

    thrust::stable_sort_by_key(h_keys.begin(), h_keys.end(), h_vals.begin());
    thrust::stable_sort_by_key(d_keys.begin(), d_keys.end(), d_vals.begin());

    ASSERT_EQ_QUIET(h_keys, d_keys);
    ASSERT_EQ_QUIET(h_vals, d_vals);
}

TEST(StableSortByKeyLargeTests, TestStableSortByKeyWithLargeKeysAndValues)
{
    SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
    
    _TestStableSortByKeyWithLargeKeysAndValues<int, 4>();
    _TestStableSortByKeyWithLargeKeysAndValues<int, 8>();
    // XXX these take too long to compile,
    _TestStableSortByKeyWithLargeKeysAndValues<int, 16>();
    _TestStableSortByKeyWithLargeKeysAndValues<int, 32>();
    _TestStableSortByKeyWithLargeKeysAndValues<int, 64>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int,  128>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int,  256>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int,  512>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int, 1024>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int, 2048>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int, 4096>();
    //    _TestStableSortByKeyWithLargeKeysAndValues<int, 8192>();
}