Integrating thrust into HPX

CMakeLists.txt

@@ -606,6 +606,20 @@ if(HPX_WITH_CUDA AND HPX_WITH_HIP)
   )
 endif()
 
+# HPX THRUST configuration
+set(HPX_WITH_THRUST_DEFAULT OFF)
+if(HPX_WITH_CUDA)
+  set(HPX_WITH_THRUST_DEFAULT ON)
+endif()
+
+hpx_option(
+  HPX_WITH_THRUST BOOL
+  "Enable support for NVIDIA Thrust integration (default: ON when CUDA is enabled, OFF otherwise)"
+  ${HPX_WITH_THRUST_DEFAULT}
+  ADVANCED
+  CATEGORY "Build Targets"
+)
+
 # ## HPX STDEXEC configuration ##
 
 hpx_option(

libs/core/CMakeLists.txt

@@ -1,4 +1,4 @@
-# Copyright (c) 2020-2021 The STE||AR-Group
+# Copyright (c) 2020-2023 The STE||AR-Group
 #
 # SPDX-License-Identifier: BSL-1.0
 # Distributed under the Boost Software License, Version 1.0. (See accompanying
@@ -79,6 +79,7 @@ set(_hpx_core_modules
     threading
     threading_base
     threadmanager
+    thrust
     timed_execution
     timing
     topology
@@ -98,4 +99,4 @@ endforeach()
 set(_hpx_core_modules
     ${_hpx_core_modules}
     PARENT_SCOPE
-)
+)

libs/core/algorithms/include/hpx/parallel/algorithms/fill.hpp

@@ -384,4 +384,4 @@ namespace hpx {
     } fill_n{};
 }    // namespace hpx
 
-#endif    // DOXYGEN
+#endif    // DOXYGEN

libs/core/execution_base/include/hpx/execution_base/any_sender.hpp

@@ -36,7 +36,7 @@ namespace hpx::detail {
     template <typename T>
     using empty_vtable_t = typename empty_vtable_type<T>::type;
 
-#if !defined(HPX_MSVC) && defined(HPX_HAVE_CXX20_TRIVIAL_VIRTUAL_DESTRUCTOR)
+    #if !defined(HPX_MSVC) && defined(HPX_HAVE_CXX20_TRIVIAL_VIRTUAL_DESTRUCTOR) && !defined(__CUDACC__) && !defined(HPX_COMPUTE_DEVICE_CODE)
     template <typename T>
     inline constexpr empty_vtable_t<T> empty_vtable{};
 

libs/core/modules.rst

@@ -81,4 +81,4 @@ Core modules
    /libs/core/topology/docs/index.rst
    /libs/core/type_support/docs/index.rst
    /libs/core/util/docs/index.rst
-   /libs/core/version/docs/index.rst
+   /libs/core/version/docs/index.rst

libs/core/thrust/CMakeLists.txt

@@ -0,0 +1,41 @@
+# Copyright (c) 2025 The STE||AR-Group
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+if(NOT (HPX_WITH_CUDA AND HPX_WITH_THRUST))
+  return()
+endif()
+
+list(APPEND CMAKE_MODULE_PATH "${CMAKE_CURRENT_SOURCE_DIR}/cmake")
+
+set(thrust_headers
+    hpx/thrust/policy.hpp
+    hpx/thrust/algorithms.hpp
+    hpx/thrust/detail/algorithm_map.hpp
+)
+
+# cmake-format: off
+set(thrust_compat_headers)
+# cmake-format: on
+
+set(thrust_sources)
+
+include(HPX_AddModule)
+add_hpx_module(
+  core thrust
+  GLOBAL_HEADER_GEN ON
+  SOURCES ${thrust_sources}
+  HEADERS ${thrust_headers}
+  COMPAT_HEADERS ${thrust_compat_headers}
+  MODULE_DEPENDENCIES
+    hpx_algorithms
+    hpx_concepts
+    hpx_execution
+    hpx_execution_base
+    hpx_functional
+    hpx_futures
+    hpx_async_cuda
+  CMAKE_SUBDIRS examples tests
+)

libs/core/thrust/docs/index.rst

@@ -0,0 +1,90 @@
+..
+    Copyright (c) 2025 The STE||AR-Group
+
+    SPDX-License-Identifier: BSL-1.0
+    Distributed under the Boost Software License, Version 1.0. (See accompanying
+    file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+.. _modules_thrust:
+
+======
+thrust
+======
+
+The thrust module integrates |hpx| parallel algorithms with NVIDIA Thrust, 
+enabling GPU acceleration using familiar |hpx| algorithm syntax.
+
+Execution Policies
+==================
+
+``hpx::thrust::thrust_host_policy``
+   CPU execution using optimized parallel implementations
+
+``hpx::thrust::thrust_device_policy``
+   Synchronous GPU execution
+
+``hpx::thrust::thrust_task_policy``
+   Asynchronous GPU execution returning |hpx| futures
+
+Policy Mappings
+===============
+
+The |hpx| thrust policies map directly to Thrust execution policies:
+
+.. code-block:: c++
+
+    // Policy mappings
+    hpx::thrust::thrust_host_policy    -> thrust::host
+    hpx::thrust::thrust_device_policy  -> thrust::device  
+    hpx::thrust::thrust_task_policy    -> thrust::par_nosync
+
+Usage
+=====
+
+Synchronous Device Execution
+-----------------------------
+
+.. code-block:: c++
+
+    #include <hpx/thrust/policy.hpp>
+    #include <hpx/thrust/algorithms.hpp>
+    #include <thrust/device_vector.h>
+
+    hpx::thrust::thrust_device_policy device{};
+    thrust::device_vector<int> d_vec(1000, 0);
+
+    hpx::fill(device, d_vec.begin(), d_vec.end(), 42);
+    int sum = hpx::reduce(device, d_vec.begin(), d_vec.end(), 0);
+
+Asynchronous Execution with Futures
+------------------------------------
+
+.. code-block:: c++
+
+    #include <hpx/thrust/policy.hpp>
+    #include <hpx/async_cuda/cuda_polling_helper.hpp>
+
+    int hpx_main() {
+        // Required for async operations
+        hpx::cuda::experimental::enable_user_polling polling_guard("default");
+
+        hpx::thrust::thrust_task_policy task{};
+        thrust::device_vector<int> d_vec(1000, 1);
+
+        auto fill_future = hpx::fill(task, d_vec.begin(), d_vec.end(), 42);
+        fill_future.get();  // Standard HPX future operations
+
+    }
+
+Build Requirements
+==================
+
+* CUDA Toolkit 12.4.0+
+* |hpx| with ``HPX_WITH_CUDA=ON``
+* Enable with: ``cmake -DHPX_WITH_CUDA=ON ...``
+
+See Also
+========
+
+* :ref:`modules_async_cuda` - CUDA integration
+* :ref:`modules_execution` - |hpx| execution policies

libs/core/thrust/examples/CMakeLists.txt

@@ -0,0 +1,34 @@
+# Copyright (c)      2025 Aditya Sapra
+#
+# SPDX-License-Identifier: BSL-1.0
+# Distributed under the Boost Software License, Version 1.0. (See accompanying
+# file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+
+if(HPX_WITH_EXAMPLES)
+  add_hpx_pseudo_target(examples.modules.thrust)
+  add_hpx_pseudo_dependencies(examples.modules examples.modules.thrust)
+  if(HPX_WITH_CUDA AND HPX_WITH_THRUST)
+    set(thrust_examples device_fill host_fill device_copy device_transform async_fill_n reverse_reduce)
+
+    foreach(example ${thrust_examples})
+      set(source ${example}.cu)
+      source_group("Source Files" FILES ${source})
+
+      add_hpx_executable(
+        ${example} INTERNAL_FLAGS
+        SOURCES ${source}
+        EXCLUDE_FROM_ALL
+        HPX_PREFIX ${HPX_BUILD_PREFIX}
+        FOLDER "Examples/Modules/Core/Thrust"
+      )
+
+      add_hpx_pseudo_dependencies(examples.modules.thrust ${example})
+    endforeach()
+  endif()
+  if(HPX_WITH_TESTS AND HPX_WITH_TESTS_EXAMPLES)
+    add_hpx_pseudo_target(tests.examples.modules.thrust)
+    add_hpx_pseudo_dependencies(
+      tests.examples.modules tests.examples.modules.thrust
+    )
+  endif()
+endif()

libs/core/thrust/examples/async_fill_n.cu

@@ -0,0 +1,35 @@
+//  Copyright (c)      2025 Aditya Sapra
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+//
+
+#include <hpx/algorithm.hpp>
+#include <hpx/execution.hpp>
+#include <hpx/init.hpp>
+
+#include <hpx/thrust/algorithms.hpp>
+#include <hpx/thrust/policy.hpp>
+
+#include <hpx/async_cuda/cuda_polling_helper.hpp>
+#include <hpx/modules/async_cuda.hpp>
+#include <thrust/device_vector.h>
+
+int hpx_main(int, char**)
+{
+    // Enable CUDA event polling on the "default" pool for HPX <-> CUDA future bridging
+    hpx::cuda::experimental::enable_user_polling polling_guard("default");
+
+    hpx::thrust::thrust_task_policy task{};
+    thrust::device_vector<int> v(1024, 0);
+    auto f = hpx::fill_n(task, v.begin(), 512, 9);
+    f.get();
+    return hpx::local::finalize();
+}
+
+int main(int argc, char** argv)
+{
+    hpx::local::init_params init_args;
+    return hpx::local::init(hpx_main, argc, argv, init_args);
+}

libs/core/thrust/examples/device_copy.cu

@@ -0,0 +1,30 @@
+//  Copyright (c)      2025 Aditya Sapra
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+//
+
+#include <hpx/algorithm.hpp>
+#include <hpx/execution.hpp>
+#include <hpx/init.hpp>
+
+#include <hpx/thrust/algorithms.hpp>
+#include <hpx/thrust/policy.hpp>
+
+#include <thrust/device_vector.h>
+
+int hpx_main(int, char**)
+{
+    hpx::thrust::thrust_device_policy device{};
+    thrust::device_vector<int> src(1024, 5);
+    thrust::device_vector<int> dst(1024, 0);
+    hpx::copy(device, src.begin(), src.end(), dst.begin());
+    return hpx::local::finalize();
+}
+
+int main(int argc, char** argv)
+{
+    hpx::local::init_params init_args;
+    return hpx::local::init(hpx_main, argc, argv, init_args);
+}

libs/core/thrust/examples/device_fill.cu

@@ -0,0 +1,29 @@
+//  Copyright (c)      2025 Aditya Sapra
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+//
+
+#include <hpx/algorithm.hpp>
+#include <hpx/execution.hpp>
+#include <hpx/init.hpp>
+
+#include <hpx/thrust/algorithms.hpp>
+#include <hpx/thrust/policy.hpp>
+
+#include <thrust/device_vector.h>
+
+int hpx_main(int, char**)
+{
+    hpx::thrust::thrust_device_policy device{};
+    thrust::device_vector<int> v(1024, 0);
+    hpx::fill(device, v.begin(), v.end(), 1);
+    return hpx::local::finalize();
+}
+
+int main(int argc, char** argv)
+{
+    hpx::local::init_params init_args;
+    return hpx::local::init(hpx_main, argc, argv, init_args);
+}

libs/core/thrust/examples/device_transform.cu

@@ -0,0 +1,33 @@
+//  Copyright (c)      2025 Aditya Sapra
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+//
+
+#include <hpx/algorithm.hpp>
+#include <hpx/execution.hpp>
+#include <hpx/init.hpp>
+
+#include <hpx/thrust/algorithms.hpp>
+#include <hpx/thrust/policy.hpp>
+
+#include <thrust/device_vector.h>
+
+int hpx_main(int, char**)
+{
+    hpx::thrust::thrust_device_policy device{};
+    thrust::device_vector<int> a(1024, 2);
+    thrust::device_vector<int> b(1024, 3);
+    thrust::device_vector<int> out(1024, 0);
+
+    hpx::transform(device, a.begin(), a.end(), b.begin(), out.begin(),
+        [] __device__(int x, int y) { return x + y; });
+    return hpx::local::finalize();
+}
+
+int main(int argc, char** argv)
+{
+    hpx::local::init_params init_args;
+    return hpx::local::init(hpx_main, argc, argv, init_args);
+}

libs/core/thrust/examples/host_fill.cu

@@ -0,0 +1,29 @@
+//  Copyright (c)      2025 Aditya Sapra
+//
+//  SPDX-License-Identifier: BSL-1.0
+//  Distributed under the Boost Software License, Version 1.0. (See accompanying
+//  file LICENSE_1_0.txt or copy at http://www.boost.org/LICENSE_1_0.txt)
+//
+
+#include <hpx/algorithm.hpp>
+#include <hpx/execution.hpp>
+#include <hpx/init.hpp>
+
+#include <hpx/thrust/algorithms.hpp>
+#include <hpx/thrust/policy.hpp>
+
+#include <vector>
+
+int hpx_main(int, char**)
+{
+    hpx::thrust::thrust_host_policy host{};
+    std::vector<int> v(1024, 0);
+    hpx::fill(host, v.begin(), v.end(), 42);
+    return hpx::local::finalize();
+}
+
+int main(int argc, char** argv)
+{
+    hpx::local::init_params init_args;
+    return hpx::local::init(hpx_main, argc, argv, init_args);
+}