NASP/cuda-samples
60
0
Fork 0
mirror of https://github.com/NVIDIA/cuda-samples.git synced 2026-04-02 15:45:41 +08:00
Code Issues Packages Projects Releases Wiki Activity

Bug 5812203: Replace thrust::make_tuple,tuple,get with the cuda::std ones

Browse Source
This commit is contained in:
Shawn Zeng 2026-01-16 17:56:50 +08:00
parent af49216f0c
commit 6c4d183ba3
6 changed files with 43 additions and 38 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
Samples/2_Concepts_and_Techniques/particles/particleSystem_cuda.cu
View File

@ -135,8 +135,8 @@ extern "C"
thrust::device_ptr<float4> d_pos4((float4 *)pos);
thrust::device_ptr<float4> 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));
}

11
Samples/2_Concepts_and_Techniques/particles/particles_kernel_impl.cuh
View File

@ -41,6 +41,9 @@
#include "thrust/iterator/zip_iterator.h"
#include "thrust/sort.h"
// for cuda::std::get
#include <cuda/std/utility>
namespace cg = cooperative_groups;
#include "helper_math.h"
#include "math_constants.h"
@ -60,8 +63,8 @@ struct integrate_functor
template <typename Tuple> __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);
}
};

23
Samples/2_Concepts_and_Techniques/segmentationTreeThrust/segmentationTree.cu
View File

@ -505,13 +505,12 @@ private:
thrust::device_ptr<uint> dMinScannedEdges = pools.uintEdges.get();
thrust::device_ptr<float> dMinScannedWeights = pools.floatEdges.get();
thrust::inclusive_scan_by_key(
dEdgesFlags,
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::make_zip_iterator(dWeights_, dEdges_),
thrust::make_zip_iterator(dMinScannedWeights, dMinScannedEdges),
thrust::greater_equal<uint>(),
thrust::minimum<thrust::tuple<float, uint>>());
thrust::minimum<cuda::std::tuple<float, uint>>());
// 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<uint>(dSuccessors),
thrust::device_ptr<uint>(dClusteredVerticesIDs))),
thrust::make_zip_iterator(
thrust::make_tuple(thrust::device_ptr<uint>(dSuccessors + verticesCount_),
thrust::device_ptr<uint>(dClusteredVerticesIDs + verticesCount_))));
thrust::sort(thrust::make_zip_iterator(thrust::device_ptr<uint>(dSuccessors),
thrust::device_ptr<uint>(dClusteredVerticesIDs)),
thrust::make_zip_iterator(thrust::device_ptr<uint>(dSuccessors + verticesCount_),
thrust::device_ptr<uint>(dClusteredVerticesIDs + verticesCount_)));
// Mark those groups.
thrust::device_ptr<uint> 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 =

12
Samples/3_CUDA_Features/cdpQuadtree/cdpQuadtree.cu
View File

@ -30,6 +30,9 @@
#include <thrust/host_vector.h>
#include <thrust/random.h>
// for cuda::std::tuple
#include <cuda/std/tuple>
namespace cg = cooperative_groups;
#include <helper_cuda.h>
@ -610,7 +613,7 @@ struct Random_generator
return a;
}
__host__ __device__ __forceinline__ thrust::tuple<float, float> operator()()
__host__ __device__ __forceinline__ cuda::std::tuple<float, float> 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<float> 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];

10
Samples/5_Domain_Specific/smokeParticles/ParticleSystem_cuda.cu
View File

@ -127,9 +127,9 @@ extern "C"
thrust::device_ptr<float4> 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<float> d_keys(keys);
thrust::device_ptr<uint> 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);

15
Samples/5_Domain_Specific/smokeParticles/particles_kernel_device.cuh
View File

@ -40,6 +40,9 @@
#include "thrust/iterator/zip_iterator.h"
#include "thrust/sort.h"
// for cuda::std::get
#include <cuda/std/utility>
cudaTextureObject_t noiseTex;
// simulation parameters
__constant__ SimParams cudaParams;
@ -65,8 +68,8 @@ struct integrate_functor
template <typename Tuple> __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 <typename Tuple> __host__ __device__ void operator()(Tuple t)
{
volatile float4 p = thrust::get<0>(t);
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;
}
};