From 6c4d183ba30202520b66e7b524f47780bb1f3c36 Mon Sep 17 00:00:00 2001 From: Shawn Zeng Date: Fri, 16 Jan 2026 17:56:50 +0800 Subject: [PATCH] Bug 5812203: Replace thrust::make_tuple,tuple,get with the cuda::std ones --- .../particles/particleSystem_cuda.cu | 4 +-- .../particles/particles_kernel_impl.cuh | 11 +++++--- .../segmentationTree.cu | 27 +++++++++---------- .../cdpQuadtree/cdpQuadtree.cu | 12 +++++---- .../smokeParticles/ParticleSystem_cuda.cu | 10 +++---- .../particles_kernel_device.cuh | 17 +++++++----- 6 files changed, 43 insertions(+), 38 deletions(-) diff --git a/Samples/2_Concepts_and_Techniques/particles/particleSystem_cuda.cu b/Samples/2_Concepts_and_Techniques/particles/particleSystem_cuda.cu index ce9e82a8..dbfa1296 100644 --- a/Samples/2_Concepts_and_Techniques/particles/particleSystem_cuda.cu +++ b/Samples/2_Concepts_and_Techniques/particles/particleSystem_cuda.cu @@ -135,8 +135,8 @@ extern "C" thrust::device_ptr d_pos4((float4 *)pos); thrust::device_ptr d_vel4((float4 *)vel); - thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(d_pos4, d_vel4)), - thrust::make_zip_iterator(thrust::make_tuple(d_pos4 + numParticles, d_vel4 + numParticles)), + thrust::for_each(thrust::make_zip_iterator(d_pos4, d_vel4), + thrust::make_zip_iterator(d_pos4 + numParticles, d_vel4 + numParticles), integrate_functor(deltaTime)); } diff --git a/Samples/2_Concepts_and_Techniques/particles/particles_kernel_impl.cuh b/Samples/2_Concepts_and_Techniques/particles/particles_kernel_impl.cuh index b855b7c4..6e47f0a5 100644 --- a/Samples/2_Concepts_and_Techniques/particles/particles_kernel_impl.cuh +++ b/Samples/2_Concepts_and_Techniques/particles/particles_kernel_impl.cuh @@ -41,6 +41,9 @@ #include "thrust/iterator/zip_iterator.h" #include "thrust/sort.h" +// for cuda::std::get +#include + namespace cg = cooperative_groups; #include "helper_math.h" #include "math_constants.h" @@ -60,8 +63,8 @@ struct integrate_functor template __device__ void operator()(Tuple t) { - volatile float4 posData = thrust::get<0>(t); - volatile float4 velData = thrust::get<1>(t); + volatile float4 posData = cuda::std::get<0>(t); + volatile float4 velData = cuda::std::get<1>(t); float3 pos = make_float3(posData.x, posData.y, posData.z); float3 vel = make_float3(velData.x, velData.y, velData.z); @@ -107,8 +110,8 @@ struct integrate_functor } // store new position and velocity - thrust::get<0>(t) = make_float4(pos, posData.w); - thrust::get<1>(t) = make_float4(vel, velData.w); + cuda::std::get<0>(t) = make_float4(pos, posData.w); + cuda::std::get<1>(t) = make_float4(vel, velData.w); } }; diff --git a/Samples/2_Concepts_and_Techniques/segmentationTreeThrust/segmentationTree.cu b/Samples/2_Concepts_and_Techniques/segmentationTreeThrust/segmentationTree.cu index b0c63605..0a265e00 100644 --- a/Samples/2_Concepts_and_Techniques/segmentationTreeThrust/segmentationTree.cu +++ b/Samples/2_Concepts_and_Techniques/segmentationTreeThrust/segmentationTree.cu @@ -505,13 +505,12 @@ private: thrust::device_ptr dMinScannedEdges = pools.uintEdges.get(); thrust::device_ptr dMinScannedWeights = pools.floatEdges.get(); - thrust::inclusive_scan_by_key( - dEdgesFlags, - dEdgesFlags + edgesCount_, - thrust::make_zip_iterator(thrust::make_tuple(dWeights_, dEdges_)), - thrust::make_zip_iterator(thrust::make_tuple(dMinScannedWeights, dMinScannedEdges)), - thrust::greater_equal(), - thrust::minimum>()); + thrust::inclusive_scan_by_key(dEdgesFlags, + dEdgesFlags + edgesCount_, + thrust::make_zip_iterator(dWeights_, dEdges_), + thrust::make_zip_iterator(dMinScannedWeights, dMinScannedEdges), + thrust::greater_equal(), + thrust::minimum>()); // To make things clear. // Let "dEdgesFlags" denote groups of edges that @@ -562,11 +561,10 @@ private: thrust::sequence(dClusteredVerticesIDs, dClusteredVerticesIDs + verticesCount_); - thrust::sort(thrust::make_zip_iterator(thrust::make_tuple(thrust::device_ptr(dSuccessors), - thrust::device_ptr(dClusteredVerticesIDs))), - thrust::make_zip_iterator( - thrust::make_tuple(thrust::device_ptr(dSuccessors + verticesCount_), - thrust::device_ptr(dClusteredVerticesIDs + verticesCount_)))); + thrust::sort(thrust::make_zip_iterator(thrust::device_ptr(dSuccessors), + thrust::device_ptr(dClusteredVerticesIDs)), + thrust::make_zip_iterator(thrust::device_ptr(dSuccessors + verticesCount_), + thrust::device_ptr(dClusteredVerticesIDs + verticesCount_))); // Mark those groups. thrust::device_ptr dVerticesFlags_ = pools.uintVertices.get(); @@ -653,9 +651,8 @@ private: // the step of the algorithm to that one. Hence we need to // preserve the structure of the original graph: neighbours and // weights should be grouped by vertex. - thrust::sort(thrust::make_zip_iterator(thrust::make_tuple(dNewStartpoints, dSurvivedEdgesIDs)), - thrust::make_zip_iterator( - thrust::make_tuple(dNewStartpoints + edgesCount_, dSurvivedEdgesIDs + edgesCount_))); + thrust::sort(thrust::make_zip_iterator(dNewStartpoints, dSurvivedEdgesIDs), + thrust::make_zip_iterator(dNewStartpoints + edgesCount_, dSurvivedEdgesIDs + edgesCount_)); // Find the group of contracted edges. uint *invalidEdgesPtr = diff --git a/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu b/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu index 60112fda..804d2550 100644 --- a/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu +++ b/Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu @@ -30,6 +30,9 @@ #include #include +// for cuda::std::tuple +#include + namespace cg = cooperative_groups; #include @@ -610,7 +613,7 @@ struct Random_generator return a; } - __host__ __device__ __forceinline__ thrust::tuple operator()() + __host__ __device__ __forceinline__ cuda::std::tuple operator()() { #ifdef __CUDA_ARCH__ unsigned seed = hash(blockIdx.x * blockDim.x + threadIdx.x + count); @@ -622,7 +625,7 @@ struct Random_generator #endif thrust::default_random_engine rng(seed); thrust::random::uniform_real_distribution distrib; - return thrust::make_tuple(distrib(rng), distrib(rng)); + return cuda::std::make_tuple(distrib(rng), distrib(rng)); } }; @@ -644,9 +647,8 @@ bool cdpQuadtree(int warp_size) // Generate random points. Random_generator rnd; - thrust::generate(thrust::make_zip_iterator(thrust::make_tuple(x_d0.begin(), y_d0.begin())), - thrust::make_zip_iterator(thrust::make_tuple(x_d0.end(), y_d0.end())), - rnd); + thrust::generate( + thrust::make_zip_iterator(x_d0.begin(), y_d0.begin()), thrust::make_zip_iterator(x_d0.end(), y_d0.end()), rnd); // Host structures to analyze the device ones. Points points_init[2]; diff --git a/Samples/5_Domain_Specific/smokeParticles/ParticleSystem_cuda.cu b/Samples/5_Domain_Specific/smokeParticles/ParticleSystem_cuda.cu index 69af1466..ab06d847 100644 --- a/Samples/5_Domain_Specific/smokeParticles/ParticleSystem_cuda.cu +++ b/Samples/5_Domain_Specific/smokeParticles/ParticleSystem_cuda.cu @@ -127,9 +127,9 @@ extern "C" thrust::device_ptr d_oldVel(oldVel); thrust::for_each( - thrust::make_zip_iterator(thrust::make_tuple(d_newPos, d_newVel, d_oldPos, d_oldVel)), - thrust::make_zip_iterator(thrust::make_tuple( - d_newPos + numParticles, d_newVel + numParticles, d_oldPos + numParticles, d_oldVel + numParticles)), + thrust::make_zip_iterator(d_newPos, d_newVel, d_oldPos, d_oldVel), + thrust::make_zip_iterator( + d_newPos + numParticles, d_newVel + numParticles, d_oldPos + numParticles, d_oldVel + numParticles), integrate_functor(deltaTime, noiseTex)); } @@ -143,8 +143,8 @@ extern "C" thrust::device_ptr d_keys(keys); thrust::device_ptr d_indices(indices); - thrust::for_each(thrust::make_zip_iterator(thrust::make_tuple(d_pos, d_keys)), - thrust::make_zip_iterator(thrust::make_tuple(d_pos + numParticles, d_keys + numParticles)), + thrust::for_each(thrust::make_zip_iterator(d_pos, d_keys), + thrust::make_zip_iterator(d_pos + numParticles, d_keys + numParticles), calcDepth_functor(sortVector)); thrust::sequence(d_indices, d_indices + numParticles); diff --git a/Samples/5_Domain_Specific/smokeParticles/particles_kernel_device.cuh b/Samples/5_Domain_Specific/smokeParticles/particles_kernel_device.cuh index 8e6e449b..f43fcd15 100644 --- a/Samples/5_Domain_Specific/smokeParticles/particles_kernel_device.cuh +++ b/Samples/5_Domain_Specific/smokeParticles/particles_kernel_device.cuh @@ -40,6 +40,9 @@ #include "thrust/iterator/zip_iterator.h" #include "thrust/sort.h" +// for cuda::std::get +#include + cudaTextureObject_t noiseTex; // simulation parameters __constant__ SimParams cudaParams; @@ -65,8 +68,8 @@ struct integrate_functor template __device__ void operator()(Tuple t) { - volatile float4 posData = thrust::get<2>(t); - volatile float4 velData = thrust::get<3>(t); + volatile float4 posData = cuda::std::get<2>(t); + volatile float4 velData = cuda::std::get<3>(t); float3 pos = make_float3(posData.x, posData.y, posData.z); float3 vel = make_float3(velData.x, velData.y, velData.z); @@ -95,8 +98,8 @@ struct integrate_functor vel *= cudaParams.globalDamping; // store new position and velocity - thrust::get<0>(t) = make_float4(pos, age); - thrust::get<1>(t) = make_float4(vel, velData.w); + cuda::std::get<0>(t) = make_float4(pos, age); + cuda::std::get<1>(t) = make_float4(vel, velData.w); } }; @@ -111,10 +114,10 @@ struct calcDepth_functor template __host__ __device__ void operator()(Tuple t) { - volatile float4 p = thrust::get<0>(t); - float key = -dot(make_float3(p.x, p.y, p.z), + volatile float4 p = cuda::std::get<0>(t); + float key = -dot(make_float3(p.x, p.y, p.z), sortVector); // project onto sort vector - thrust::get<1>(t) = key; + cuda::std::get<1>(t) = key; } };