Merge branch 'staging'

StreamHPC · Aug 9, 2022 · 85931df · 85931df
2 parents e934a3a + 767898e
commit 85931df
Show file tree

Hide file tree

Showing 12 changed files with 200 additions and 235 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -2,6 +2,10 @@
 
 Full documentation for rocThrust is available at [https://rocthrust.readthedocs.io/en/latest/](https://rocthrust.readthedocs.io/en/latest/)
 
+## (Unreleased) rocThrust 2.16.0 for ROCm 5.3
+### Changed
+- rocThrust functionality dependent on device malloc works is functional as ROCm 5.2 reneabled device malloc. Device launched `thrust::sort` and `thrust::sort_by_key` are available for use.
+
 ## rocThrust 2.15.0 for ROCm 5.2
 ### Added
 - Packages for tests and benchmark executable on all supported OSes using CPack.

diff --git a/README.md b/README.md
@@ -8,7 +8,6 @@ Thrust is a parallel algorithm library. This library has been ported to [HIP](ht
 
 ### Software
 
-* Git
 * CMake (3.5.1 or later)
 * AMD [ROCm](https://rocm.github.io/install.html) platform (1.8.0 or later)
   * Including [HipCC](https://github.com/ROCm-Developer-Tools/HIP) compiler, which must be

diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake
@@ -12,12 +12,6 @@
 # For downloading, building, and installing required dependencies
 include(cmake/DownloadProject.cmake)
 
-# GIT
-find_package(Git REQUIRED)
-if (NOT Git_FOUND)
-  message(FATAL_ERROR "Please ensure Git is installed on the system")
-endif()
-
 # rocPRIM (https://github.com/ROCmSoftwarePlatform/rocPRIM)
 if(NOT DOWNLOAD_ROCPRIM)
   find_package(rocprim QUIET)

diff --git a/test/test_partition.cpp b/test/test_partition.cpp
@@ -1906,8 +1906,6 @@ TEST(PartitionTests,TestPartitionCopyDevice)
       }
 }
 
-//TODO: Tests fails with error message "Memory access fault by GPU node-1 (Agent handle: 0x1ade7d0) on address 0x7fac25a00000. Reason: Page not present or supervisor privilege."
-
 // __global__
 // THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
 // void StablePartitionKernel(int const N, int* array)

diff --git a/test/test_sort.cpp b/test/test_sort.cpp
@@ -437,45 +437,47 @@ TEST(SortTests, TestSortBoolDescending)
     }
 }
 
+#ifndef _WIN32
 //TODO: refactor this test into a different set of tests
-// __global__
-// THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
-// void SortKernel(int const N, int* array)
-// {
-//     if(threadIdx.x == 0)
-//     {
-//         thrust::device_ptr<int> begin(array);
-//         thrust::device_ptr<int> end(array + N);
-//         thrust::sort(thrust::hip::par, begin, end);
-//     }
-// }
-//
-// TEST(SortTests, TestSortDevice)
-// {
-//     SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
-//
-//     for(auto size : get_sizes() )
-//     {
-//         SCOPED_TRACE(testing::Message() << "with size= " << size);
-//
-//         for(auto seed : get_seeds())
-//         {
-//             SCOPED_TRACE(testing::Message() << "with seed= " << seed);
-//
-//             thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
-//
-//             thrust::device_vector<int> d_data = h_data;
-//
-//             thrust::sort(h_data.begin(), h_data.end());
-//             hipLaunchKernelGGL(SortKernel,
-//                                dim3(1, 1, 1),
-//                                dim3(128, 1, 1),
-//                                0,
-//                                0,
-//                                size,
-//                                thrust::raw_pointer_cast(&d_data[0]));
-//
-//             ASSERT_EQ(h_data, d_data);
-//         }
-//     }
-// }
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void SortKernel(int const N, int* array)
+{
+    if(threadIdx.x == 0)
+    {
+        thrust::device_ptr<int> begin(array);
+        thrust::device_ptr<int> end(array + N);
+        thrust::sort(thrust::hip::par, begin, end);
+    }
+}
+
+TEST(SortTests, TestSortDevice)
+{
+    SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+    for(auto size : get_sizes() )
+    {
+        SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+        for(auto seed : get_seeds())
+        {
+            SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+            thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
+
+            thrust::device_vector<int> d_data = h_data;
+
+            thrust::sort(h_data.begin(), h_data.end());
+            hipLaunchKernelGGL(SortKernel,
+                               dim3(1, 1, 1),
+                               dim3(128, 1, 1),
+                               0,
+                               0,
+                               size,
+                               thrust::raw_pointer_cast(&d_data[0]));
+
+            ASSERT_EQ(h_data, d_data);
+        }
+    }
+}
+#endif
diff --git a/test/test_sort_by_key.cpp b/test/test_sort_by_key.cpp
@@ -256,57 +256,57 @@ TEST(SortByKeyTests, TestSortByKeyBoolDescending)
     }
 }
 
-//TODO: Tests fails with error message "Memory access fault by GPU node-1 (Agent handle: 0x1ade7d0) on address 0x7fac25a00000. Reason: Page not present or supervisor privilege."
-
-// __global__
-// THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
-// void SortByKeyKernel(int const N, int* keys, short* values)
-// {
-//     if(threadIdx.x == 0)
-//     {
-//         thrust::device_ptr<int>   keys_begin(keys);
-//         thrust::device_ptr<int>   keys_end(keys + N);
-//         thrust::device_ptr<short> val(values);
-//         thrust::sort_by_key(thrust::hip::par, keys_begin, keys_end, val);
-//     }
-// }
-//
-// TEST(SortByKeyTests, TestSortByKeyDevice)
-// {
-//     SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
-//
-//     for (auto size: get_sizes() )
-//     {
-//         SCOPED_TRACE(testing::Message() << "with size= " << size);
-//
-//         for(auto seed : get_seeds())
-//         {
-//             SCOPED_TRACE(testing::Message() << "with seed= " << seed);
-//
-//             thrust::host_vector<int> h_keys = get_random_data<int>(size, 0, size, seed);
-//
-//             thrust::host_vector<short> h_values
-//                 = get_random_data<short>(size,
-//                                          std::numeric_limits<short>::min(),
-//                                          std::numeric_limits<short>::max(),
-//                                          seed);
-//
-//             thrust::device_vector<int>   d_keys   = h_keys;
-//             thrust::device_vector<short> d_values = h_values;
-//
-//             thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
-//             hipLaunchKernelGGL(SortByKeyKernel,
-//                                dim3(1, 1, 1),
-//                                dim3(128, 1, 1),
-//                                0,
-//                                0,
-//                                size,
-//                                thrust::raw_pointer_cast(&d_keys[0]),
-//                                thrust::raw_pointer_cast(&d_values[0]));
-//
-//             ASSERT_EQ(h_keys, d_keys);
-//             // Only keys are compared here, the sequential stable_merge_sort that's used in
-//             // CUDA and HIP don't generate the correct value sorting
-//         }
-//     }
-// }
+#ifndef _WIN32
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void SortByKeyKernel(int const N, int* keys, short* values)
+{
+    if(threadIdx.x == 0)
+    {
+        thrust::device_ptr<int>   keys_begin(keys);
+        thrust::device_ptr<int>   keys_end(keys + N);
+        thrust::device_ptr<short> val(values);
+        thrust::sort_by_key(thrust::hip::par, keys_begin, keys_end, val);
+    }
+}
+
+TEST(SortByKeyTests, TestSortByKeyDevice)
+{
+    SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+    for (auto size: get_sizes() )
+    {
+        SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+        for(auto seed : get_seeds())
+        {
+            SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+            thrust::host_vector<int> h_keys = get_random_data<int>(size, 0, size, seed);
+
+            thrust::host_vector<short> h_values
+                = get_random_data<short>(size,
+                                         std::numeric_limits<short>::min(),
+                                         std::numeric_limits<short>::max(),
+                                         seed);
+
+            thrust::device_vector<int>   d_keys   = h_keys;
+            thrust::device_vector<short> d_values = h_values;
+
+            thrust::sort_by_key(h_keys.begin(), h_keys.end(), h_values.begin());
+            hipLaunchKernelGGL(SortByKeyKernel,
+                               dim3(1, 1, 1),
+                               dim3(128, 1, 1),
+                               0,
+                               0,
+                               size,
+                               thrust::raw_pointer_cast(&d_keys[0]),
+                               thrust::raw_pointer_cast(&d_values[0]));
+
+            ASSERT_EQ(h_keys, d_keys);
+            // Only keys are compared here, the sequential stable_merge_sort that's used in
+            // CUDA and HIP don't generate the correct value sorting
+        }
+    }
+}
+#endif
diff --git a/test/test_stable_sort.cpp b/test/test_stable_sort.cpp
@@ -189,45 +189,46 @@ TYPED_TEST(StableSortVectorTests, TestStableSortWithIndirection)
     ASSERT_EQ(data[6], T(2));
 }
 
+#ifndef _WIN32
+__global__
+THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
+void StableSortKernel(int const N, int* array)
+{
+    if(threadIdx.x == 0)
+    {
+        thrust::device_ptr<int> begin(array);
+        thrust::device_ptr<int> end(array + N);
+        thrust::stable_sort(thrust::hip::par, begin, end);
+    }
+}
 
-// __global__
-// THRUST_HIP_LAUNCH_BOUNDS_DEFAULT
-// void StableSortKernel(int const N, int* array)
-// {
-//     if(threadIdx.x == 0)
-//     {
-//         thrust::device_ptr<int> begin(array);
-//         thrust::device_ptr<int> end(array + N);
-//         thrust::stable_sort(thrust::hip::par, begin, end);
-//     }
-// }
-//
-// TEST(StableSortTests, TestStableSortDevice)
-// {
-//     SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
-//
-//     for(auto size : get_sizes() )
-//     {
-//         SCOPED_TRACE(testing::Message() << "with size= " << size);
-//
-//         for(auto seed : get_seeds())
-//         {
-//             SCOPED_TRACE(testing::Message() << "with seed= " << seed);
-//
-//             thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
-//
-//             thrust::device_vector<int> d_data = h_data;
-//
-//             thrust::stable_sort(h_data.begin(), h_data.end());
-//             hipLaunchKernelGGL(StableSortKernel,
-//                                dim3(1, 1, 1),
-//                                dim3(128, 1, 1),
-//                                0,
-//                                0,
-//                                size,
-//                                thrust::raw_pointer_cast(&d_data[0]));
-// 
-//             ASSERT_EQ(h_data, d_data);
-//         }
-//     }
-// }
+TEST(StableSortTests, TestStableSortDevice)
+{
+    SCOPED_TRACE(testing::Message() << "with device_id= " << test::set_device_from_ctest());
+
+    for(auto size : get_sizes() )
+    {
+        SCOPED_TRACE(testing::Message() << "with size= " << size);
+
+        for(auto seed : get_seeds())
+        {
+            SCOPED_TRACE(testing::Message() << "with seed= " << seed);
+
+            thrust::host_vector<int> h_data = get_random_data<int>(size, 0, size, seed);
+
+            thrust::device_vector<int> d_data = h_data;
+
+            thrust::stable_sort(h_data.begin(), h_data.end());
+            hipLaunchKernelGGL(StableSortKernel,
+                               dim3(1, 1, 1),
+                               dim3(128, 1, 1),
+                               0,
+                               0,
+                               size,
+                               thrust::raw_pointer_cast(&d_data[0]));
+
+            ASSERT_EQ(h_data, d_data);
+        }
+    }
+}
+#endif