Refactored Memory Wrappers tests, fixed uninitialized_copy
MODIFY
test/test_device_delete.cpp
test/test_device_delete.cpp
@@ -66,7 +66,7 @@ TEST(DeviceDelete, TestDeviceDeleteDestructorInvocation)
ASSERT_EQ(false, destructor_flag[0]);
- // TODO
+ // TODO: Known failure
//thrust::device_delete(foo_ptr);
//ASSERT_EQ(true, destructor_flag[0]);
}

MODIFY
test/test_device_ptr.cpp
test/test_device_ptr.cpp
#include <gtest/gtest.h>
// Thrust
-#include <thrust/memory.h>
-#include <thrust/transform.h>
-// STREAMHPC TODO replace <thrust/detail/seq.h> with <thrust/execution_policy.h>
-#include <thrust/detail/seq.h>
-#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
+#include <thrust/device_ptr.h>
// HIP API
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
@@ -56,6 +52,34 @@ public:
using input_type = typename Params::input_type;
};
+typedef ::testing::Types<
+ Params<thrust::host_vector<short>>,
+ Params<thrust::host_vector<int>>,
+ Params<thrust::host_vector<long long>>,
+ Params<thrust::host_vector<unsigned short>>,
+ Params<thrust::host_vector<unsigned int>>,
+ Params<thrust::host_vector<unsigned long long>>,
+ Params<thrust::host_vector<float>>,
+ Params<thrust::host_vector<double>>,
+ Params<thrust::device_vector<short>>,
+ Params<thrust::device_vector<int>>,
+ Params<thrust::device_vector<long long>>,
+ Params<thrust::device_vector<unsigned short>>,
+ Params<thrust::device_vector<unsigned int>>,
+ Params<thrust::device_vector<unsigned long long>>,
+ Params<thrust::device_vector<float>>,
+ Params<thrust::device_vector<double>>
+> DevicePtrTestsParams;
+
+TYPED_TEST_CASE(DevicePtrTests, DevicePtrTestsParams);
+
+template<class Params>
+class DevicePtrPrimitiveTests : public ::testing::Test
+{
+public:
+ using input_type = typename Params::input_type;
+};
+
typedef ::testing::Types<
Params<short>,
Params<int>,
@@ -65,9 +89,9 @@ typedef ::testing::Types<
Params<unsigned long long>,
Params<float>,
Params<double>
-> DevicePtrTestsParams;
+> DevicePtrPrimitiveTestsParams;
-TYPED_TEST_CASE(DevicePtrTests, DevicePtrTestsParams);
+TYPED_TEST_CASE(DevicePtrPrimitiveTests, DevicePtrPrimitiveTestsParams);
template <typename T>
struct mark_processed_functor
@@ -78,18 +102,6 @@ struct mark_processed_functor
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
-TYPED_TEST(DevicePtrTests, MakeDevicePointer)
-{
- using T = typename TestFixture::input_type;
-
- T *raw_ptr = 0;
- thrust::device_ptr<T> p0 = thrust::device_pointer_cast(raw_ptr);
-
- ASSERT_EQ(thrust::raw_pointer_cast(p0), raw_ptr);
- thrust::device_ptr<T> p1 = thrust::device_pointer_cast(p0);
- ASSERT_EQ(p0, p1);
-}
-
TEST(DevicePtrTests,TestDevicePointerManipulation)
{
thrust::device_vector<int> data(5);
@@ -135,10 +147,24 @@ TEST(DevicePtrTests,TestDevicePointerManipulation)
ASSERT_EQ(end - begin, 5);
}
-TYPED_TEST(DevicePtrTests,TestRawPointerCast)
+TYPED_TEST(DevicePtrPrimitiveTests, MakeDevicePointer)
{
using T = typename TestFixture::input_type;
- thrust::device_vector<T> vec(3);
+
+ T *raw_ptr = 0;
+ thrust::device_ptr<T> p0 = thrust::device_pointer_cast(raw_ptr);
+
+ ASSERT_EQ(thrust::raw_pointer_cast(p0), raw_ptr);
+ thrust::device_ptr<T> p1 = thrust::device_pointer_cast(p0);
+ ASSERT_EQ(p0, p1);
+}
+
+TYPED_TEST(DevicePtrTests,TestRawPointerCast)
+{
+ using Vector = typename TestFixture::input_type;
+ using T = typename Vector::value_type;
+
+ Vector vec(3);
T * first;
T * last;
@@ -152,7 +178,7 @@ TYPED_TEST(DevicePtrTests,TestRawPointerCast)
ASSERT_EQ(last - first, 2);
}
-TYPED_TEST(DevicePtrTests, TestDevicePointerValue)
+TYPED_TEST(DevicePtrPrimitiveTests, TestDevicePointerValue)
{
using T = typename TestFixture::input_type;

MODIFY
test/test_device_reference.cpp
test/test_device_reference.cpp
#include <gtest/gtest.h>
// Thrust
-#include <thrust/memory.h>
-#include <thrust/transform.h>
-// STREAMHPC TODO replace <thrust/detail/seq.h> with <thrust/execution_policy.h>
-#include <thrust/detail/seq.h>
-#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
+#include <thrust/device_reference.h>
// HIP API
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
@@ -90,8 +86,6 @@ typedef ::testing::Types<
TYPED_TEST_CASE(DeviceReferenceIntegerTests, DeviceReferenceIntegerTestsParams);
-
-
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
TYPED_TEST(DeviceReferenceTests, TestDeviceReferenceConstructorFromDeviceReference)
@@ -143,7 +137,6 @@ TYPED_TEST(DeviceReferenceTests, TestDeviceReferenceConstructorFromDevicePointer
ASSERT_EQ(v[0], ref);
}
-// TODO: No device code available for function: _ZN7rocprim6detail16transform_kernelILj256ELj16ERKjN6thrust10device_ptrIjEES6_NS4_8identityIjEEEEvT2_mT3_T4_
TEST(DeviceReferenceTests, TestDeviceReferenceAssignmentFromDeviceReference)
{
// test same types

MODIFY
test/test_memory.cpp
test/test_memory.cpp
// Thrust
#include <thrust/memory.h>
+#include <thrust/sort.h>
+#include <thrust/memory.h>
+#include <thrust/pair.h>
+#include <thrust/fill.h>
+#include <thrust/logical.h>
+#include <thrust/sequence.h>
+#include <thrust/reverse.h>
+#include <thrust/execution_policy.h>
// HIP API
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
TEST(HipThrustMemory, VoidMalloc)
{
- const size_t n = 9001;
+ const size_t size = 9001;
thrust::device_system_tag dev_tag;
using pointer = thrust::pointer<int, thrust::device_system_tag>;
// Malloc on device
- auto void_ptr = thrust::malloc(dev_tag, sizeof(int) * n);
+ auto void_ptr = thrust::malloc(dev_tag, sizeof(int) * size);
pointer ptr = pointer(static_cast<int*>(void_ptr.get()));
// Free
thrust::free(dev_tag, ptr);
@@ -52,11 +60,11 @@ TEST(HipThrustMemory, VoidMalloc)
TEST(HipThrustMemory, TypeMalloc)
{
- const size_t n = 9001;
+ const size_t size = 9001;
thrust::device_system_tag dev_tag;
// Malloc on device
- auto ptr = thrust::malloc<int>(dev_tag, sizeof(int) * n);
+ auto ptr = thrust::malloc<int>(dev_tag, sizeof(int) * size);
// Free
thrust::free(dev_tag, ptr);
}
@@ -64,16 +72,230 @@ TEST(HipThrustMemory, TypeMalloc)
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
TEST(HipThrustMemory, MallocUseMemory)
{
- const size_t n = 1024;
+ const size_t size = 1024;
thrust::device_system_tag dev_tag;
// Malloc on device
- auto ptr = thrust::malloc<int>(dev_tag, sizeof(int) * n);
+ auto ptr = thrust::malloc<int>(dev_tag, sizeof(int) * size);
// Try allocated memory with HIP function
- HIP_CHECK(hipMemset(ptr.get(), 0, n * sizeof(int)));
+ HIP_CHECK(hipMemset(ptr.get(), 0, size * sizeof(int)));
// Free
thrust::free(dev_tag, ptr);
}
#endif // THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC
+
+// Define a new system class, as the my_system one is already used with a thrust::sort template definition
+// that calls back into sort.cu
+class my_memory_system : public thrust::device_execution_policy<my_memory_system>
+{
+ public:
+ my_memory_system(int)
+ : correctly_dispatched(false),
+ num_copies(0)
+ {}
+
+ my_memory_system(const my_memory_system &other)
+ : correctly_dispatched(false),
+ num_copies(other.num_copies + 1)
+ {}
+
+ void validate_dispatch()
+ {
+ correctly_dispatched = (num_copies == 0);
+ }
+
+ bool is_valid()
+ {
+ return correctly_dispatched;
+ }
+
+ private:
+ bool correctly_dispatched;
+
+ // count the number of copies so that we can validate
+ // that dispatch does not introduce any
+ unsigned int num_copies;
+
+
+ // disallow default construction
+ my_memory_system();
+};
+
+template<typename T1, typename T2>
+bool are_same(const T1 &, const T2 &)
+{
+ return false;
+}
+
+
+template<typename T>
+bool are_same(const T &, const T &)
+{
+ return true;
+}
+
+TEST(MemoryTests, TestSelectSystemDifferentTypes)
+{
+ using thrust::system::detail::generic::select_system;
+
+ my_memory_system my_sys(0);
+ thrust::device_system_tag device_sys;
+
+ // select_system(my_system, device_system_tag) should return device_system_tag (the minimum tag)
+ bool is_device_system_tag = are_same(device_sys, select_system(my_sys, device_sys));
+ ASSERT_EQ(true, is_device_system_tag);
+
+ // select_system(device_system_tag, my_tag) should return device_system_tag (the minimum tag)
+ is_device_system_tag = are_same(device_sys, select_system(device_sys, my_sys));
+ ASSERT_EQ(true, is_device_system_tag);
+}
+
+TEST(MemoryTests, TestSelectSystemSameTypes)
+{
+ using thrust::system::detail::generic::select_system;
+
+ my_memory_system my_sys(0);
+ thrust::device_system_tag device_sys;
+ thrust::host_system_tag host_sys;
+
+ // select_system(host_system_tag, host_system_tag) should return host_system_tag
+ bool is_host_system_tag = are_same(host_sys, select_system(host_sys, host_sys));
+ ASSERT_EQ(true, is_host_system_tag);
+
+ // select_system(device_system_tag, device_system_tag) should return device_system_tag
+ bool is_device_system_tag = are_same(device_sys, select_system(device_sys, device_sys));
+ ASSERT_EQ(true, is_device_system_tag);
+
+ // select_system(my_system, my_system) should return my_system
+ bool is_my_system = are_same(my_sys, select_system(my_sys, my_sys));
+ ASSERT_EQ(true, is_my_system);
+}
+
+TEST(MemoryTests, TestGetTemporaryBuffer)
+{
+ const size_t size = 9001;
+
+ thrust::device_system_tag dev_tag;
+ typedef thrust::pointer<int, thrust::device_system_tag> pointer;
+ thrust::pair<pointer, std::ptrdiff_t> ptr_and_sz = thrust::get_temporary_buffer<int>(dev_tag, size);
+
+ ASSERT_EQ(ptr_and_sz.second, size);
+
+ const int ref_val = 13;
+ thrust::device_vector<int> ref(size, ref_val);
+
+ thrust::fill_n(ptr_and_sz.first, size, ref_val);
+
+ ASSERT_EQ(true, thrust::all_of(ptr_and_sz.first, ptr_and_sz.first + size, thrust::placeholders::_1 == ref_val));
+
+ thrust::return_temporary_buffer(dev_tag, ptr_and_sz.first);
+}
+
+TEST(MemoryTests, TestMalloc)
+{
+ const size_t size = 9001;
+
+ thrust::device_system_tag dev_tag;
+ typedef thrust::pointer<int, thrust::device_system_tag> pointer;
+ pointer ptr = pointer(static_cast<int*>(thrust::malloc(dev_tag, sizeof(int) * size).get()));
+
+ const int ref_val = 13;
+ thrust::device_vector<int> ref(size, ref_val);
+
+ thrust::fill_n(ptr, size, ref_val);
+
+ ASSERT_EQ(true, thrust::all_of(ptr, ptr + size, thrust::placeholders::_1 == ref_val));
+
+ thrust::free(dev_tag, ptr);
+}
+
+thrust::pointer<void,my_memory_system>
+ malloc(my_memory_system &system, std::size_t)
+{
+ system.validate_dispatch();
+
+ return thrust::pointer<void,my_memory_system>();
+}
+
+TEST(MemoryTests, TestMallocDispatchExplicit)
+{
+ const size_t size = 0;
+
+ my_memory_system sys(0);
+ thrust::malloc(sys, size);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename Pointer>
+void free(my_memory_system &system, Pointer)
+{
+ system.validate_dispatch();
+}
+
+TEST(MemoryTests, TestFreeDispatchExplicit)
+{
+ thrust::pointer<my_memory_system,void> ptr;
+
+ my_memory_system sys(0);
+ thrust::free(sys, ptr);
+
+ ASSERT_EQ(true, sys.is_valid());
+}
+
+template<typename T>
+ thrust::pair<thrust::pointer<T,my_memory_system>, std::ptrdiff_t>
+ get_temporary_buffer(my_memory_system &system, std::ptrdiff_t size)
+{
+ system.validate_dispatch();
+
+ thrust::device_system_tag device_sys;
+ thrust::pair<thrust::pointer<T, thrust::device_system_tag>, std::ptrdiff_t> result = thrust::get_temporary_buffer<T>(device_sys, size);
+ return thrust::make_pair(thrust::pointer<T,my_memory_system>(result.first.get()), result.second);
+}
+
+TEST(MemoryTests, TestGetTemporaryBufferDispatchImplicit)
+{
+ const size_t size = 9001;
+
+ my_memory_system sys(0);
+ typedef thrust::pointer<int, thrust::device_system_tag> pointer;
+ thrust::pair<pointer, std::ptrdiff_t> ptr_and_sz = thrust::get_temporary_buffer<int>(sys, size);
+
+ ASSERT_EQ(ptr_and_sz.second, size);
+ ASSERT_EQ(true, sys.is_valid());
+
+ const int ref_val = 13;
+ thrust::device_vector<int> ref(size, ref_val);
+
+ thrust::fill_n(ptr_and_sz.first, size, ref_val);
+
+ ASSERT_EQ(true, thrust::all_of(ptr_and_sz.first, ptr_and_sz.first + size, thrust::placeholders::_1 == ref_val));
+
+ thrust::return_temporary_buffer(sys, ptr_and_sz.first);
+}
+
+TEST(MemoryTests, TestGetTemporaryBufferDispatchExplicit)
+{
+ if(are_same(thrust::device_system_tag(), thrust::system::cpp::tag()))
+ {
+ // XXX cpp uses the internal scalar backend, which currently elides user tags
+ // KNOWN_FAILURE;
+ }
+ else
+ {
+ thrust::device_vector<int> vec(9001);
+
+ thrust::sequence(vec.begin(), vec.end());
+ thrust::reverse(vec.begin(), vec.end());
+
+ // call something we know will invoke get_temporary_buffer
+ my_memory_system sys(0);
+ thrust::sort(sys, vec.begin(), vec.end());
+
+ ASSERT_EQ(true, thrust::is_sorted(vec.begin(), vec.end()));
+ ASSERT_EQ(true, sys.is_valid());
+ }
+}

MODIFY
test/test_uninitialized_copy.cpp
test/test_uninitialized_copy.cpp
@@ -234,8 +234,7 @@ struct CopyConstructTest
bool copy_constructed_on_device;
};
-// TODO: Compare with cuda tests and find the cause of differences
-/*TEST(UninitializedCopyTests, TestUninitializedCopyNonPODDevice)
+TEST(UninitializedCopyTests, TestUninitializedCopyNonPODDevice)
{
using T = CopyConstructTest;
@@ -246,7 +245,6 @@ struct CopyConstructTest
ASSERT_EQ(false, x.copy_constructed_on_host);
x = v1[0];
- // TODO: Check why it has te be true?
ASSERT_EQ(true, x.copy_constructed_on_device);
ASSERT_EQ(false, x.copy_constructed_on_host);
@@ -268,7 +266,6 @@ TEST(UninitializedCopyTests, TestUninitializedCopyNNonPODDevice)
ASSERT_EQ(false, x.copy_constructed_on_host);
x = v1[0];
- // TODO: Check why it has te be true?
ASSERT_EQ(true, x.copy_constructed_on_device);
ASSERT_EQ(false, x.copy_constructed_on_host);
@@ -277,7 +274,7 @@ TEST(UninitializedCopyTests, TestUninitializedCopyNNonPODDevice)
x = v2[0];
ASSERT_EQ(true, x.copy_constructed_on_device);
ASSERT_EQ(false, x.copy_constructed_on_host);
-}*/
+}
TEST(UninitializedCopyTests, TestUninitializedCopyNonPODHost)
{

MODIFY
test/test_vector.cpp
test/test_vector.cpp
@@ -148,10 +148,7 @@ TYPED_TEST(VectorTests, TestVectorElementAssignment)
using Vector = typename TestFixture::input_type;
using T = typename Vector::value_type;
- // TODO: No device code available for function: _ZN7rocprim6detail16transform_kernelILj256ELj16ERKjN6thrust10device_ptrIjEES6_NS4_8identityIjEEEEvT2_mT3_T4_
- // Vector v(3);
- thrust::host_vector<T> v(3);
-
+ Vector v(3);
v[0] = T(0); v[1] = T(1); v[2] = T(2);
ASSERT_EQ(v[0], T(0));
@@ -164,9 +161,7 @@ TYPED_TEST(VectorTests, TestVectorElementAssignment)
ASSERT_EQ(v[1], T(11));
ASSERT_EQ(v[2], T(12));
- // TODO: No device code available for function: _ZN7rocprim6detail16transform_kernelILj256ELj16ERKjN6thrust10device_ptrIjEES6_NS4_8identityIjEEEEvT2_mT3_T4_
- // Vector w(3);
- thrust::host_vector<T> w(3);
+ Vector w(3);
w[0] = v[0];
w[1] = v[1];
w[2] = v[2];
@@ -260,9 +255,7 @@ TYPED_TEST(VectorTests, TestVectorAssignFromBiDirectionalIterator)
stl_list.push_back(T(1));
stl_list.push_back(T(2));
- // TODO: No device code available for function: _ZN7rocprim6detail16transform_kernelILj256ELj16ERKjN6thrust10device_ptrIjEES6_NS4_8identityIjEEEEvT2_mT3_T4_
- // Vector v;
- thrust::host_vector<int> v;
+ Vector v;
v.assign(stl_list.begin(), stl_list.end());
ASSERT_EQ(v.size(), 3);

MODIFY
test/test_vector_insert.cpp
test/test_vector_insert.cpp
#include <gtest/gtest.h>
// Thrust
-#include <thrust/memory.h>
-#include <thrust/transform.h>
-// STREAMHPC TODO replace <thrust/detail/seq.h> with <thrust/execution_policy.h>
#include <thrust/sequence.h>
#include <thrust/device_malloc_allocator.h>
-#include <thrust/device_vector.h>
-#include <thrust/host_vector.h>
// HIP API
#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_HCC

MODIFY
test/test_vector_manipulation.cpp
test/test_vector_manipulation.cpp
@@ -87,19 +87,9 @@ TYPED_TEST(VectorManipulationTests, TestVectorManipulation)
const std::vector<size_t> sizes = get_sizes();
for(auto size : sizes)
{
- thrust::host_vector<T> src;
- if (std::is_floating_point<T>::value)
- {
- src = get_random_data<T>(size, (T)-1000, (T)+1000);
- }
- else
- {
- src = get_random_data<T>(
- size,
+ thrust::host_vector<T> src = get_random_data<T>(size,
std::numeric_limits<T>::min(),
- std::numeric_limits<T>::max()
- );
- }
+ std::numeric_limits<T>::max());
ASSERT_EQ(src.size(), size);

ADD
thrust/system/hip/detail/logical.h
+/*
+ * Copyright 2008-2013 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.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special version of this algorithm
+

MODIFY
thrust/system/hip/detail/uninitialized_copy.h
thrust/system/hip/detail/uninitialized_copy.h
@@ -59,13 +59,7 @@ namespace __uninitialized_copy {
InputType const &in = raw_reference_cast(input[idx]);
OutputType & out = raw_reference_cast(output[idx]);
-// STREAMHPC Check if this is a problem for HCC too
-#if defined(__HCC__) && defined(__clang__)
- // XXX unsafe, but clang is seemngly unable to call in-place new
- out = in;
-#else
::new (static_cast<void *>(&out)) OutputType(in);
-#endif
}
}; // struct functor

Refactored Memory Wrappers tests, fixed uninitialized_copy