StreamHPC 2023-10-30 (#338)

* fix: Restore compatibility with latest rocPRIM

rocPRIM changed the API of `lookback_scan_state`, update usage to match.

* fix: doxygen warnings

---------

Co-authored-by: Gergely Meszaros <gergely@streamhpc.com>

Author: Naraenda

thrust/mr/new.h

@@ -1,5 +1,6 @@
 /*
  *  Copyright 2018 NVIDIA Corporation
+ *  Modifications Copyright 2023 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.
@@ -39,6 +40,13 @@ namespace mr
 class new_delete_resource final : public memory_resource<>
 {
 public:
+    /*! Allocates memory of size at least \p bytes and alignment at least \p alignment.
+     *
+     *  \param bytes size, in bytes, that is requested from this allocation
+     *  \param alignment alignment that is requested from this allocation
+     *  \throws thrust::bad_alloc when no memory with requested size and alignment can be allocated.
+     *  \return A pointer to void to the newly allocated memory.
+     */
     void * do_allocate(std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override
     {
 #if defined(__cpp_aligned_new)
@@ -60,6 +68,14 @@ class new_delete_resource final : public memory_resource<>
 #endif
     }
 
+    /*! Deallocates memory pointed to by \p p.
+     *
+     *  \param p pointer to be deallocated
+     *  \param bytes the size of the allocation. This must be equivalent to the value of \p bytes that
+     *      was passed to the allocation function that returned \p p.
+     *  \param alignment the size of the allocation. This must be equivalent to the value of \p alignment
+     *      that was passed to the allocation function that returned \p p.
+     */
     void do_deallocate(void * p, std::size_t bytes, std::size_t alignment = THRUST_MR_DEFAULT_ALIGNMENT) override
     {
 #if defined(__cpp_aligned_new)

thrust/system/hip/detail/set_operations.h

@@ -918,8 +918,12 @@ namespace __set_operations
         const unsigned int number_of_blocks = (keys_total + items_per_block - 1) / items_per_block;
 
         // Calculate required temporary storage
-        size_t scan_state_bytes
-            = ::rocprim::detail::align_size(block_state_type::get_storage_size(number_of_blocks));
+        size_t scan_state_bytes;
+        status = block_state_type::get_storage_size(number_of_blocks, stream, scan_state_bytes);
+        if (status != hipSuccess) {
+            return status;
+        }
+        scan_state_bytes = ::rocprim::detail::align_size(scan_state_bytes);
         size_t ordered_block_id_bytes
             = ::rocprim::detail::align_size(ordered_block_id_type::get_storage_size());
         size_t partition_storage_bytes = (number_of_blocks + 1) * sizeof(pair<Size, Size>);
@@ -943,7 +947,11 @@ namespace __set_operations
 
         auto ptr = reinterpret_cast<char*>(temporary_storage);
         // Create and initialize lookback_scan_state obj
-        auto blocks_state = block_state_type::create(ptr, number_of_blocks);
+        block_state_type blocks_state; 
+        status = block_state_type::create(blocks_state, ptr, number_of_blocks, stream);
+        if (status != hipSuccess) {
+            return status;
+        }
         ptr += scan_state_bytes;
         // Create and initialize ordered_block_id obj
         auto ordered_bid