Adding initial HMM test for device_ptr.
MODIFY
CMakeLists.txt
CMakeLists.txt
@@ -43,7 +43,7 @@ endif()
# Use target ID syntax if supported for AMDGPU_TARGETS
if(TARGET_ID_SUPPORT)
- set(AMDGPU_TARGETS gfx900:xnack-;gfx906:xnack-;gfx908:xnack- CACHE STRING "List of specific machine types for library to target")
+ set(AMDGPU_TARGETS gfx900:xnack-;gfx906:xnack-;gfx908:xnack-;gfx908:xnack+ CACHE STRING "List of specific machine types for library to target")
else()
set(AMDGPU_TARGETS gfx900;gfx906;gfx908 CACHE STRING "List of specific machine types for library to target")
endif()

MODIFY
test/test_device_ptr.cpp
test/test_device_ptr.cpp
#include <thrust/device_ptr.h>
#include <thrust/device_vector.h>
+#include <cstdlib>
#include "test_header.hpp"
+#define HIP_CHECK(condition) \
+{ \
+ hipError_t error = condition; \
+ if(error != hipSuccess){ \
+ std::cout << "HIP error: " << error << " line: " << __LINE__ << std::endl; \
+ exit(error); \
+ } \
+}
+
TESTS_DEFINE(DevicePtrTests, FullTestsParams);
TESTS_DEFINE(DevicePtrPrimitiveTests, NumericalTestsParams);
@@ -33,6 +43,22 @@ struct mark_processed_functor
}
};
+bool supports_hmm()
+{
+ hipDeviceProp_t device_prop;
+ int device_id;
+ HIP_CHECK(hipGetDevice(&device_id));
+ HIP_CHECK(hipGetDeviceProperties(&device_prop, device_id));
+ if (device_prop.managedMemory == 1) return true;
+
+ return false;
+}
+
+bool use_hmm()
+{
+ return std::getenv("ROCTHRUST_USE_HMM");
+}
+
TEST(DevicePtrTests, TestDevicePointerManipulation)
{
SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
@@ -80,6 +106,58 @@ TEST(DevicePtrTests, TestDevicePointerManipulation)
ASSERT_EQ(end - begin, 5);
}
+TEST(DevicePtrTests, TestDevicePointerManipulationHmm)
+{
+ SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+ if (!(supports_hmm() && use_hmm())) return;
+
+ int* data;
+ HIP_CHECK(hipMallocManaged((void**)&data, 5 * sizeof(int)));
+
+ thrust::device_ptr<int> begin(&data[0]);
+ thrust::device_ptr<int> end(&data[0] + 5);
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin++;
+ begin--;
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin += 1;
+ begin -= 1;
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin = begin + (int)1;
+ begin = begin - (int)1;
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin = begin + (unsigned int)1;
+ begin = begin - (unsigned int)1;
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin = begin + (size_t)1;
+ begin = begin - (size_t)1;
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin = begin + (ptrdiff_t)1;
+ begin = begin - (ptrdiff_t)1;
+
+ ASSERT_EQ(end - begin, 5);
+
+ begin = begin + (thrust::device_ptr<int>::difference_type)1;
+ begin = begin - (thrust::device_ptr<int>::difference_type)1;
+
+ ASSERT_EQ(end - begin, 5);
+
+ HIP_CHECK(hipFree(data));
+}
+
TYPED_TEST(DevicePtrPrimitiveTests, MakeDevicePointer)
{
using T = typename TestFixture::input_type;

Adding initial HMM test for device_ptr.