From 7ca23d4e594864b816b3335ddf0c741e0988d84e Mon Sep 17 00:00:00 2001 From: Kenneth Moreland Date: Thu, 28 May 2026 09:34:55 -0400 Subject: [PATCH] Fix error about ambiguous thrust namespace There are a few places in Viskores that have to modify the internal behavior of Thrust to, for example, properly identify a reference object as a read/write reference. However, recent versions of thrust sometimes caused a compiler error about the `thrust::detail` namespace being ambiguous. This is because Thrust has its special ABI that declares things in anonymous spaces to avoid conflicts. When Viskores touches internal components of Thrust, it now uses the `THRUST_NAMESPACE_BEGIN`/`END` macros to define the namespace so it follows how things are defined by Thrust. --- docs/changelog/thrust-ambiguous-namespace.md | 11 ++++++++ .../cuda/internal/IteratorFromArrayPortal.h | 6 ++--- viskores/exec/cuda/internal/ThrustPatches.h | 27 ++++++++++++------- .../exec/cuda/internal/WrappedOperators.h | 5 ++-- viskores/filter/contour/worklet/Clip.h | 5 ++-- 5 files changed, 35 insertions(+), 19 deletions(-) create mode 100644 docs/changelog/thrust-ambiguous-namespace.md diff --git a/docs/changelog/thrust-ambiguous-namespace.md b/docs/changelog/thrust-ambiguous-namespace.md new file mode 100644 index 0000000000..fdec1b3eac --- /dev/null +++ b/docs/changelog/thrust-ambiguous-namespace.md @@ -0,0 +1,11 @@ +## Fixed error about ambiguous thrust namespace + +There are a few places in Viskores that have to modify the internal behavior of +Thrust to, for example, properly identify a reference object as a read/write +reference. However, recent versions of thrust sometimes caused a compiler error +about the `thrust::detail` namespace being ambiguous. This is because Thrust has +its special ABI that declares things in anonymous spaces to avoid conflicts. + +When Viskores touches internal components of Thrust, it now uses the +`THRUST_NAMESPACE_BEGIN`/`END` macros to define the namespace so it follows how +things are defined by Thrust. diff --git a/viskores/exec/cuda/internal/IteratorFromArrayPortal.h b/viskores/exec/cuda/internal/IteratorFromArrayPortal.h index e050a86e9c..d4d136bd1d 100644 --- a/viskores/exec/cuda/internal/IteratorFromArrayPortal.h +++ b/viskores/exec/cuda/internal/IteratorFromArrayPortal.h @@ -20,6 +20,7 @@ #include #include +#include #include #include @@ -128,8 +129,7 @@ class IteratorFromArrayPortal //must be treated as references. So do to do that properly we need to specialize //is_non_const_reference to state an ArrayPortalValueReference by value is valid //for writing -namespace thrust -{ +VISKORES_THRUST_NAMESPACE_BEGIN namespace detail { @@ -142,6 +142,6 @@ struct is_non_const_reference> { }; } -} +VISKORES_THRUST_NAMESPACE_END #endif //viskores_exec_cuda_internal_IteratorFromArrayPortal_h diff --git a/viskores/exec/cuda/internal/ThrustPatches.h b/viskores/exec/cuda/internal/ThrustPatches.h index 67e5c369cf..4bafd7cbad 100644 --- a/viskores/exec/cuda/internal/ThrustPatches.h +++ b/viskores/exec/cuda/internal/ThrustPatches.h @@ -25,12 +25,21 @@ // Needed so we can conditionally include components #include +#if THRUST_VERSION >= 101301 +#define VISKORES_THRUST_NAMESPACE_BEGIN THRUST_NAMESPACE_BEGIN +#define VISKORES_THRUST_NAMESPACE_END THRUST_NAMESPACE_END +#else +#define VISKORES_THRUST_NAMESPACE_BEGIN \ + namespace thrust \ + { +#define VISKORES_THRUST_NAMESPACE_END } +#endif + #if THRUST_VERSION >= 100900 && THRUST_VERSION < 100906 //So for thrust 1.9.0+ ( CUDA 9.X+ ) the aligned_reinterpret_cast has a bug //where it is not marked as __host__device__. To fix this we add a new //overload for void* with the correct markup (which is what everyone calls). -namespace thrust -{ +VISKORES_THRUST_NAMESPACE_BEGIN namespace detail { //just in-case somebody has this fix also for primitive types @@ -179,7 +188,7 @@ ALIGN_RE_PAIR(viskores::Int64, viskores::Float64); #undef ALIGN_RE_VEC #undef ALIGN_RE_PAIR } -} +VISKORES_THRUST_NAMESPACE_END #endif //THRUST_VERSION >= 100900 && THRUST_VERSION < 100906 #if (THRUST_VERSION >= 100904) && (THRUST_VERSION < 100909) @@ -192,8 +201,7 @@ VISKORES_THIRDPARTY_PRE_INCLUDE #include #include VISKORES_THIRDPARTY_POST_INCLUDE -namespace thrust -{ +VISKORES_THRUST_NAMESPACE_BEGIN namespace mr { @@ -245,7 +253,7 @@ class stateless_resource_allocator __host__ __device__ ~stateless_resource_allocator() {} }; } -} +VISKORES_THRUST_NAMESPACE_END #endif //(THRUST_VERSION >= 100904) && (THRUST_VERSION < 100909) @@ -273,8 +281,7 @@ class WrappedBinaryOperator; } } //namespace viskores::exec::cuda::internal -namespace thrust -{ +VISKORES_THRUST_NAMESPACE_BEGIN namespace system { namespace cuda @@ -340,8 +347,8 @@ destructive_accumulate_n(ConcurrentGroup& g, } } } -} //namespace thrust::system::cuda::detail -#endif //THRUST_VERSION < 100900 +VISKORES_THRUST_NAMESPACE_END //namespace thrust::system::cuda::detail +#endif //THRUST_VERSION < 100900 #endif //CUDA enabled diff --git a/viskores/exec/cuda/internal/WrappedOperators.h b/viskores/exec/cuda/internal/WrappedOperators.h index fded81addd..60d6c19c3d 100644 --- a/viskores/exec/cuda/internal/WrappedOperators.h +++ b/viskores/exec/cuda/internal/WrappedOperators.h @@ -199,8 +199,7 @@ struct WrappedBinaryPredicate } } //namespace viskores::exec::cuda::internal -namespace thrust -{ +VISKORES_THRUST_NAMESPACE_BEGIN namespace detail { // @@ -223,6 +222,6 @@ struct is_commutative : public true_type { }; } -} // namespace thrust::detail +VISKORES_THRUST_NAMESPACE_END // namespace thrust::detail #endif #endif // viskores_m_worklet_Clip_h