Support THRUST_DEVICE_SYSTEM_CPP host backend

.github/workflows/presubmit.yml

@@ -36,6 +36,41 @@ jobs:
           module load oneapi/2025.0.4 llvm/22.0-native HIP/chipStar/main level-zero/dgpu HIP/rocPRIM/2026.04.08
           module list
           rm -rf ~/.cache/chipStar build
+          # hipcc self-locates and injects -I<its install>/include ahead of
+          # all project -I flags, so tests would compile against the thrust
+          # previously installed into the shared chipStar prefix instead of
+          # this checkout. Build a job-local hardlink copy of the install
+          # with this PR's thrust overlaid and use its hipcc.
+          HIPROOT=$(dirname "$(dirname "$(realpath "$(which hipcc)")")")
+          rm -rf hip-local
+          # If the hardlink copy fails (cross-filesystem) it leaves a partial
+          # tree behind; remove it before the plain-copy fallback or cp -a
+          # nests the install inside it.
+          cp -al "$HIPROOT" hip-local 2>/dev/null || { rm -rf hip-local; cp -a "$HIPROOT" hip-local; }
+          # Unlink before overlaying: writing through hardlinks would edit
+          # the shared install. The generated rocthrust_version.hpp is
+          # resolved from the build tree's -I instead.
+          rm -rf hip-local/include/thrust
+          mkdir -p hip-local/include/thrust
+          cp -r thrust/. hip-local/include/thrust/
+          # hipcc derives its injected -I<...>/include from the HIP_PATH
+          # recorded inside share/.hipInfo (which may differ from the
+          # install's realpath), so rewrite the whole line (sed -i replaces
+          # via rename; the shared install's file is untouched).
+          sed -i "s|^HIP_PATH=.*|HIP_PATH=$PWD/hip-local|" hip-local/share/.hipInfo
+          # hipcc's getHipPath() prefers the HIP_PATH env var (set by the
+          # module) over .hipInfo — repoint it too.
+          export HIP_PATH="$PWD/hip-local"
+          export PATH="$PWD/hip-local/bin:$PATH"
+          hash -r
+          # Fail fast if <thrust/...> still resolves outside the job-local
+          # copy — otherwise the tests silently build against the thrust
+          # bundled in the shared chipStar install instead of this checkout.
+          printf '#include <thrust/version.h>\n' > /tmp/thrust-probe.hip
+          hipcc -x hip -E /tmp/thrust-probe.hip > /tmp/thrust-probe.out
+          grep 'thrust/version.h' /tmp/thrust-probe.out | sort -u | head -2 || true
+          grep -q "hip-local/include/thrust/version.h" /tmp/thrust-probe.out \
+            || { echo "ERROR: thrust resolves outside hip-local"; exit 1; }
           cmake -S . -B build \
             -DCMAKE_CXX_COMPILER="$(which hipcc)" \
             -DCMAKE_BUILD_TYPE=Release \

test/CMakeLists.txt

@@ -120,6 +120,22 @@ endfunction()
 # Tests
 # ****************************************************************************
 
+# Host-only (CPP device system) build test — reproducer for
+# https://github.com/CHIP-SPV/chipStar/issues/1363. Compiled with
+# THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CPP so everything runs on the
+# host; no GPU or GTest needed.
+add_executable(thrust_cpp_backend.hip test_thrust_cpp_backend.cpp)
+target_compile_definitions(thrust_cpp_backend.hip
+    PRIVATE
+        THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CPP
+)
+target_link_libraries(thrust_cpp_backend.hip
+    PRIVATE
+        rocthrust
+        roc::rocprim_hip
+)
+add_relative_test("thrust_cpp_backend" thrust_cpp_backend.hip)
+
 add_rocthrust_test("adjacent_difference")
 add_rocthrust_test("advance")
 add_rocthrust_test("allocator")

test/test_thrust_cpp_backend.cpp

@@ -0,0 +1,66 @@
+// Reproducer for https://github.com/CHIP-SPV/chipStar/issues/1363
+//
+// Building any Thrust code with -DTHRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CPP
+// (the host-only fallback backend) must compile and run correctly. Before the
+// fix this failed to compile with:
+//   thrust/detail/type_traits.h: 'invoke_result_t' had no CPP-system branch
+//   thrust/detail/type_traits/result_of_adaptable_function.h: no 'invoke_result_t'
+//   thrust/system/detail/sequential/trivial_copy.h: undeclared 'NV_IS_HOST'
+//
+// This test is compiled with THRUST_DEVICE_SYSTEM=THRUST_DEVICE_SYSTEM_CPP
+// (see test/CMakeLists.txt) and runs entirely on the host.
+
+#include <thrust/device_vector.h>
+#include <thrust/host_vector.h>
+#include <thrust/reduce.h>
+#include <thrust/sort.h>
+#include <thrust/transform.h>
+#include <thrust/transform_reduce.h>
+
+#include <cassert>
+#include <cstdio>
+
+struct square
+{
+  __host__ __device__ int operator()(int x) const { return x * x; }
+};
+
+int main()
+{
+  const int n = 100;
+
+  thrust::host_vector<int> h(n);
+  for (int i = 0; i < n; ++i)
+    h[i] = i + 1;
+
+  // With the CPP device system, device_vector lives in host memory.
+  thrust::device_vector<int> d = h;
+
+  // reduce: 1 + 2 + ... + 100 = 5050
+  const int sum = thrust::reduce(d.begin(), d.end(), 0);
+  assert(sum == 5050);
+
+  // transform: element-wise square
+  thrust::device_vector<int> squared(n);
+  thrust::transform(d.begin(), d.end(), squared.begin(), square());
+  const thrust::host_vector<int> h_squared = squared;
+  for (int i = 0; i < n; ++i)
+    assert(h_squared[i] == (i + 1) * (i + 1));
+
+  // transform_reduce: sum of squares = n(n+1)(2n+1)/6 = 338350
+  const int sum_sq = thrust::transform_reduce(d.begin(), d.end(), square(), 0,
+                                              thrust::plus<int>());
+  assert(sum_sq == 338350);
+
+  // sort a reversed copy
+  thrust::device_vector<int> rev(n);
+  for (int i = 0; i < n; ++i)
+    rev[i] = n - i;
+  thrust::sort(rev.begin(), rev.end());
+  const thrust::host_vector<int> h_sorted = rev;
+  for (int i = 0; i < n; ++i)
+    assert(h_sorted[i] == i + 1);
+
+  std::printf("PASS\n");
+  return 0;
+}

thrust/detail/nv_target.h

@@ -27,4 +27,10 @@
 #  include <thrust/system/hip/detail/nv/target.h>
 #elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
 #  include <nv/target>
+#else
+// Host-only device systems (CPP, OMP, TBB) have no separate device pass of
+// their own. The bundled shim below is backend-agnostic: it only keys on
+// __HIP_DEVICE_COMPILE__ (never defined for pure host compilation), so
+// NV_IF_TARGET(NV_IS_HOST, ...) statically selects the host branch.
+#  include <thrust/system/hip/detail/nv/target.h>
 #endif

thrust/detail/type_traits.h

@@ -731,6 +731,12 @@ using invoke_result_t =
 #else // 2017+
   ::cuda::std::invoke_result_t<Invokable, Args...>;
 #endif
+#else // host-only device systems (CPP, OMP, TBB)
+#if THRUST_CPP_DIALECT < 2017
+  typename ::std::result_of<Invokable(Args...)>::type;
+#else // 2017+
+  ::std::invoke_result_t<Invokable, Args...>;
+#endif
 #endif
 
 template <class F, class... Us>