From 49355f1e869395dbf299539bfec307ed1e812d9e Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 22 Jul 2026 10:11:19 +0300 Subject: [PATCH 1/6] tests: add reproducer for CPP device system build (#1363) --- test/CMakeLists.txt | 16 ++++++++ test/test_thrust_cpp_backend.cpp | 66 ++++++++++++++++++++++++++++++++ 2 files changed, 82 insertions(+) create mode 100644 test/test_thrust_cpp_backend.cpp diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index ee25f6f53..735b8ccd8 100644 --- a/test/CMakeLists.txt +++ b/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") diff --git a/test/test_thrust_cpp_backend.cpp b/test/test_thrust_cpp_backend.cpp new file mode 100644 index 000000000..f8c69a017 --- /dev/null +++ b/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 +#include +#include +#include +#include +#include + +#include +#include + +struct square +{ + __host__ __device__ int operator()(int x) const { return x * x; } +}; + +int main() +{ + const int n = 100; + + thrust::host_vector 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 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 squared(n); + thrust::transform(d.begin(), d.end(), squared.begin(), square()); + const thrust::host_vector 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()); + assert(sum_sq == 338350); + + // sort a reversed copy + thrust::device_vector rev(n); + for (int i = 0; i < n; ++i) + rev[i] = n - i; + thrust::sort(rev.begin(), rev.end()); + const thrust::host_vector h_sorted = rev; + for (int i = 0; i < n; ++i) + assert(h_sorted[i] == i + 1); + + std::printf("PASS\n"); + return 0; +} From e9bb25de11f98a0cad58e9d1408ef97f9aecedfd Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Wed, 22 Jul 2026 10:15:20 +0300 Subject: [PATCH 2/6] fix: support THRUST_DEVICE_SYSTEM_CPP host backend (#1363) --- thrust/detail/nv_target.h | 6 ++++++ thrust/detail/type_traits.h | 6 ++++++ 2 files changed, 12 insertions(+) diff --git a/thrust/detail/nv_target.h b/thrust/detail/nv_target.h index 1f93a096b..5cf7777db 100644 --- a/thrust/detail/nv_target.h +++ b/thrust/detail/nv_target.h @@ -27,4 +27,10 @@ # include #elif THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA # include +#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 #endif diff --git a/thrust/detail/type_traits.h b/thrust/detail/type_traits.h index 61532ad44..5803d677a 100644 --- a/thrust/detail/type_traits.h +++ b/thrust/detail/type_traits.h @@ -731,6 +731,12 @@ using invoke_result_t = #else // 2017+ ::cuda::std::invoke_result_t; #endif +#else // host-only device systems (CPP, OMP, TBB) +#if THRUST_CPP_DIALECT < 2017 + typename ::std::result_of::type; +#else // 2017+ + ::std::invoke_result_t; +#endif #endif template From d397fcf62ecb0da2956debe659cf5ae74cd8dc9e Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 23 Jul 2026 12:47:03 +0300 Subject: [PATCH 3/6] ci: compile tests against this checkout's thrust via job-local chipStar overlay --- .github/workflows/presubmit.yml | 16 ++++++++++++++++ 1 file changed, 16 insertions(+) diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index 3dfb3d2f9..d2942df90 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -36,6 +36,22 @@ 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/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 + cp -al "$HIPROOT" hip-local 2>/dev/null || 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/ + export PATH="$PWD/hip-local/bin:$PATH" + hash -r cmake -S . -B build \ -DCMAKE_CXX_COMPILER="$(which hipcc)" \ -DCMAKE_BUILD_TYPE=Release \ From 1a79d711725213ebcba3fe2fe5cd9232f6d5b6e5 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 23 Jul 2026 14:09:25 +0300 Subject: [PATCH 4/6] ci: repoint copied .hipInfo HIP_PATH at the job-local install --- .github/workflows/presubmit.yml | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index d2942df90..d367daaa3 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -50,6 +50,11 @@ jobs: 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, not from its own location — + # repoint it at the job-local copy (sed -i replaces via rename, + # so the shared install's file is untouched). + sed -i "s|$HIPROOT|$PWD/hip-local|g" hip-local/share/.hipInfo export PATH="$PWD/hip-local/bin:$PATH" hash -r cmake -S . -B build \ From 33ea20ea5777bb4f4cfc1391232f351976010526 Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 23 Jul 2026 14:53:17 +0300 Subject: [PATCH 5/6] ci: fix cross-fs copy fallback, rewrite .hipInfo HIP_PATH line, add resolution probe --- .github/workflows/presubmit.yml | 21 ++++++++++++++++----- 1 file changed, 16 insertions(+), 5 deletions(-) diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index d367daaa3..e8f5e316d 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -43,7 +43,10 @@ jobs: # with this PR's thrust overlaid and use its hipcc. HIPROOT=$(dirname "$(dirname "$(realpath "$(which hipcc)")")") rm -rf hip-local - cp -al "$HIPROOT" hip-local 2>/dev/null || cp -a "$HIPROOT" 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. @@ -51,12 +54,20 @@ jobs: 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, not from its own location — - # repoint it at the job-local copy (sed -i replaces via rename, - # so the shared install's file is untouched). - sed -i "s|$HIPROOT|$PWD/hip-local|g" hip-local/share/.hipInfo + # 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 export PATH="$PWD/hip-local/bin:$PATH" hash -r + # Fail fast if 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 \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 \ From ef5526dd8819aebef4fea974bac2a2e455fc132c Mon Sep 17 00:00:00 2001 From: Paulius Velesko Date: Thu, 23 Jul 2026 21:38:14 +0300 Subject: [PATCH 6/6] ci: repoint HIP_PATH env var at the job-local install (hipcc prefers it over .hipInfo) --- .github/workflows/presubmit.yml | 3 +++ 1 file changed, 3 insertions(+) diff --git a/.github/workflows/presubmit.yml b/.github/workflows/presubmit.yml index e8f5e316d..f32bbac0f 100644 --- a/.github/workflows/presubmit.yml +++ b/.github/workflows/presubmit.yml @@ -58,6 +58,9 @@ jobs: # 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 still resolves outside the job-local