thrust 相关问题

为什么transform_reduce结果与transform&reduce结果不同？

我正在测试一些推力代码，发现transform_reduce给出的计算结果略有不同，这让我完全困惑。 这是一个测试示例代码：（计算 sum(exp(x))） 原来是

CUB 按键减少

Thrust 有算法： 按键排序 按键减少 这对于我的问题来说可以很好地协同工作。我想尝试使用 CUB 来更好地控制内存和流以及与我的交互......

.cu 和 .cpp 文件之间共享 CUDA Thrust 对象

似乎最近对 Thrust（CUDA Toolkit 的一部分）进行了更改，以使用编译时 CUDA 架构来标记 host_vector 类名。当我...

CUDA的thrust::inclusive_scan()有'init'参数吗？

根据CUDA的Thrust库文档，thrust::inclusive_scan()有4个参数： OutputIterator推力::inclusive_scan（首先是InputIterator，

是否有可能克服thrust::zip_iterator中迭代器的最大数量？

我在工作中使用 Thrust 来完成一些任务，并发现在构造 zip_iterator 时似乎存在最大数量的迭代器。 例如 #包括 我在工作中使用 Thrust 来完成一些任务，并发现在构建 zip_iterator 时似乎存在最大数量的迭代器。 例如 #include <thrust/iterator/zip_iterator.h> #include <thrust/device_vector.h> int main() { thrust::device_vector<int> A(10),B(10),C(10); auto zitor = thrust::make_zip_iterator(A.begin(),A.begin(), B.begin(),B.begin(), B.begin(),B.begin(), B.begin(),B.begin(), C.begin(),C.begin()); } 这段代码编译成功。但是如果我在函数参数中再添加一个迭代器，就会出现错误： multizip.cu(8): error: no instance of overloaded function "thrust::make_zip_iterator" matches the argument list argument types are: (thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>) 1 error detected in the compilation of "multizip.cu". 它最多接受十个迭代器，还是我误解了什么？如果是这样，有什么解决办法吗？ using Itor = thrust::device_vector<int>::iterator; using Zitor = thrust::zip_iterator<thrust::tuple< Itor,Itor,Itor,Itor,Itor,Itor, Itor,Itor,Itor,Itor,Itor>>; 这段代码也不起作用，除非我删除一个Itor，错误信息是： multizip.cu(17): error: too many arguments for class template "thrust::tuple" 所以我相信thrust::tuple中的迭代器数量确实有最大限制为10。可以克服这个限制吗？ 您遇到的最接近的问题是，直到最近，Thrust 的模板限制为 thrust::zip_iterator 中的 10 个项目（迭代器）。修复/解决该问题有以下三种选项： 最近（似乎从 CUDA 12.3 到 CUDA 12.4），Thrust zip_iterator 设计已更改，允许 超过 10 个迭代器 构建 zip_iterator。 因此，一种选择是将您的 CUDA 工具包版本更新到 12.4.1 或更高版本。 另一个相关选项是将更新版本的 CCCL* 与您当前的 CUDA 工具包结合使用。 CCCL 的最新版本为此提供了“定义的兼容性路径”。 我不会在这里给出完整的秘诀，但简而言之，假设您位于定义的兼容性路径内，您将克隆计算机上当前的 CCCL 存储库，然后使用 -I* 将您的 CUDA 编译器指向该存储库*。 Thrust 是一个仅包含模板/标头的库，因此不需要单独的编译步骤或任何其他安装步骤。 根据您的需求，保持在旧 Thrust 版本的 10 个迭代器限制之内，您可以创建 zip_iterator 的 zip_iterators。 这里就是一个例子。 这可用于获得超过 10 个迭代器，尽管是嵌套排列。 *：CUDA C++ Core L库于 2023 年引入，由于它们的交互/重叠，将 Thrust、CUB 和 libcu++/libcudacxx 包含在单个存储库中。 原始的 Thrust 存储库 停止更新，并且 不 包含现代化的 zip_iterator。**：请勿使用 -isystem，因为这会获得比 CUDA 工具包打包的 CCCL 标头更低的优先级。

是否有可能克服 `thrust::zip_iterator` 中迭代器的最大数量？

我在工作中使用 Thrust 来完成一些任务，并发现在构造 zip_iterator 时似乎存在最大数量的迭代器。 例如 #包括 我在工作中使用 Thrust 来完成一些任务，并发现在构造 zip_iterator 时似乎存在最大数量的迭代器。 例如 #include <thrust/iterator/zip_iterator.h> #include <thrust/device_vector.h> int main() { thrust::device_vector<int> A(10),B(10),C(10); auto zitor = thrust::make_zip_iterator(A.begin(),A.begin(), B.begin(),B.begin(), B.begin(),B.begin(), B.begin(),B.begin(), C.begin(),C.begin()); } 这段代码编译成功。但是如果我在初始化列表中再添加一个迭代器，就会出现错误： multizip.cu(8): error: no instance of overloaded function "thrust::make_zip_iterator" matches the argument list argument types are: (thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>, thrust::detail::normal_iterator<thrust::device_ptr<int>>) 1 error detected in the compilation of "multizip.cu". 它最多接受十个迭代器，还是我误解了什么？如果是这样，有什么解决方法吗？ using Itor = thrust::device_vector<int>::iterator; using Zitor = thrust::zip_iterator<thrust::tuple< Itor,Itor,Itor,Itor,Itor,Itor, Itor,Itor,Itor,Itor,Itor>>; 这段代码也不起作用，除非我删除一个Itor，错误信息是： multizip.cu(17): error: too many arguments for class template "thrust::tuple" 所以我相信thrust::tuple中的迭代器数量确实有最大限制为10。我可以克服这个限制吗？ 迭代器的数量在编译时由模板类型决定。这就是您收到最后一个错误的原因。我认为编译器对您创建的迭代器数量存在问题。这称为重载函数（因此第一个错误说“没有重载函数的实例”）。 我可能建议您首先尝试使用两个函数，将迭代器创建分成两个独特的函数，看看这是否是问题所在。 请参阅此处：https://forums.developer.nvidia.com/t/thrust-zip-iterator-with-任意-number-of-iterators/292942。 https://forums.developer.nvidia.com/t/combining-thrust-zip-iterator-transform-iterator-counting-iterator-for-modified-summed-area-table/56154.

如何使用推力和CUDA流将内存从主机异步复制到设备

我想使用推力将内存从主机复制到设备，如下所示 推力::主机向量 h_vec(1 << 28); thrust::device_vector d_vec(1 << 28); thrust::c...

CUDA Thrust Kernels 可以在多个流上并行运行吗？

我正在尝试在不同的 CUDA 流上并行启动推力::填充两个不同的设备向量。然而，当我查看 NSight Systems 中的内核启动时，它们似乎是串行的......

尝试在需要在主机代码中查询其返回类型的上下文中使用扩展的 __device__ lambda

我收到编译器错误 static_assert 失败：“尝试在需要在主机代码中查询其返回类型的上下文中使用扩展的 __device__ lambda。使用命名函数对象，a

从 std::deque 复制到推力 device_vector

我正在尝试制作一个示例代码，从 std::deque 复制到推力::device_vector，但编译器警告不允许从 __host__ __device__ 函数调用 __host__ 函数（我...

CUDA 排序 Z 轴 3D 数组 C++/Thrust

我正在寻找沿 z 轴对大型 3D 数组进行排序。 示例数组为 X x Y x Z (1000x1000x5) 我想沿 z 轴排序，因此我会对沿 z 轴的 5 个元素执行 1000x1000 排序。 编辑

插入用户编写的内核

我是 Thrust 的新手。我看到所有 Thrust 演示文稿和示例仅显示主机代码。 我想知道我是否可以将 device_vector 传递给我自己的内核？如何？ 如果是的话，操作是什么...

使用推力来处理 CUDA 类中的向量？

我有一个关于 C++ 类的推力的适用性的问题。 我正在尝试实现一个类对象，该对象接收顶点的 (x,y,z) 坐标作为 ver1、ver2 和 ver3。然后，分配...

在CUDA工具包中包含的Thrust库中找不到thrust/universal_vector.h

我目前正在使用 Thrust 在 GPU 和 CPU 之间传输数据。但是当我在代码中包含 并使用 CMake 配置项目时，“致命错误：没有这样的...

Thrust：使用device_ptr时如何获取copy_if函数复制的元素数量

我正在使用 Thrust 库的 Thrust::copy_if 函数以及计数迭代器来获取数组中非零元素的索引。我还需要获取复制元素的数量。 我...

推力收集/过滤

我想要做的是在向量上创建一个过滤器，以便它删除未通过谓词测试的元素；但不太确定我该怎么做。 我再次评估输入向量中的每个元素...

CUDA强制指令执行顺序

我正在尝试将一些数据操作从CPU传输到GPU（CUDA），但有一小部分需要指令以特定顺序运行。原则上我可以做前几页...

CUDA 二阶递归推力 Include_scan

我试图了解如何并行递归计算。连续地，计算采用以下形式： 对于 (int i = 2; i 我试图了解如何并行递归计算。连续地，计算采用以下形式： for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } 对于i-1索引，我之前的问题有一个解决方案：CUDA强制指令执行顺序 我想修改它以使用i-2，但我不明白如何将相同的过程应用于二阶计算。应该可以使用 thrust::inclusive_scan 函数，但我不知道如何使用。有谁知道解决办法吗 从上一个问题/答案中继续，我们将注意力转移到 Blelloch 的参考论文中的方程 1.11。我们观察到您的问题表述： for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } 如果我们设置 m=2， 似乎与方程 1.11 中的情况相匹配，在这种情况下，我们还可以观察到，对于您的公式，所有 ai,1 均为零（并且，如前所述，所有 ai,2 均为零） k). 根据该论文中的方程 1.12，我们的状态变量 si 现在变成了一个二元组： si = |xi xi-1| 记下这些事情，我们观察方程 1.13 的“正确性”： si = |xi-1 xi-2| 。 |0 1，k 0| + |bi 0| 重写： si,1 = xi = k*xi-2 + bi si,2 = xi-1 = xi-1 （在我看来，其他答案让你在这一点上。这种认识，即result.data[0] = right + k * left.data[1];足以进行串行扫描，但不适用于并行扫描。同样明显的是，函子/扫描操作不具有关联性.) 我们现在需要提出一个二元运算符bop，它是（1.7）中定义对这种情况的扩展。参考之前方程1.7中的定义，我们在1.13中的处理基础上将其扩展如下： Ci = |Ai , Bi| 地点： Ai = |0 1, k 0| 和： Bi = |bi 0| 然后我们有： Ci bop Cj = | Ai。 Aj，Bi。 Aj + Bj | 这将成为我们函子/扫描运算符的公式。我们需要始终携带 6 个标量“状态”量：2 个用于 B 向量，4 个用于 A 矩阵。 接下来就是上面的实现： $ cat t1930.cu #include <iostream> #include <thrust/device_vector.h> #include <thrust/host_vector.h> #include <thrust/scan.h> #include <thrust/copy.h> #include <thrust/iterator/zip_iterator.h> #include <thrust/iterator/constant_iterator.h> #include <cstdlib> #include <cstdio> template <typename T> void cpufunction(T *result, T *oldArray, size_t size, T k){ for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } } struct scan_op // as per blelloch (1.7) { template <typename T1, typename T2> __host__ __device__ T1 operator()(const T1 &t1, const T2 &t2){ T1 ret; thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<2>(t2) + thrust::get<1>(t1)*thrust::get<4>(t2)+thrust::get<0>(t2); thrust::get<1>(ret) = thrust::get<0>(t1)*thrust::get<3>(t2) + thrust::get<1>(t1)*thrust::get<5>(t2)+thrust::get<1>(t2); thrust::get<2>(ret) = thrust::get<2>(t1)*thrust::get<2>(t2) + thrust::get<3>(t1)*thrust::get<4>(t2); thrust::get<3>(ret) = thrust::get<2>(t1)*thrust::get<3>(t2) + thrust::get<3>(t1)*thrust::get<5>(t2); thrust::get<4>(ret) = thrust::get<4>(t1)*thrust::get<2>(t2) + thrust::get<5>(t1)*thrust::get<4>(t2); thrust::get<5>(ret) = thrust::get<4>(t1)*thrust::get<3>(t2) + thrust::get<5>(t1)*thrust::get<5>(t2); return ret; } }; typedef float mt; const size_t ds = 512; const mt k = 1.01; const int snip = 10; int main(){ mt *b1 = new mt[ds]; // b as in blelloch (1.5) mt *cr = new mt[ds]; // cpu result for (int i = 0; i < ds; i++) { b1[i] = rand()/(float)RAND_MAX;} cr[0] = b1[0]; cr[1] = b1[1]; cpufunction(cr, b1, ds, k); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; thrust::device_vector<mt> db(b1, b1+ds); auto b0 = thrust::constant_iterator<mt>(0); auto a0 = thrust::constant_iterator<mt>(0); auto a1 = thrust::constant_iterator<mt>(1); auto a2 = thrust::constant_iterator<mt>(k); auto a3 = thrust::constant_iterator<mt>(0); thrust::device_vector<mt> dx1(ds); thrust::device_vector<mt> dx0(ds); thrust::device_vector<mt> dy0(ds); thrust::device_vector<mt> dy1(ds); thrust::device_vector<mt> dy2(ds); thrust::device_vector<mt> dy3(ds); auto my_i_zip = thrust::make_zip_iterator(thrust::make_tuple(db.begin(), b0, a0, a1, a2, a3)); auto my_o_zip = thrust::make_zip_iterator(thrust::make_tuple(dx1.begin(), dx0.begin(), dy0.begin(), dy1.begin(), dy2.begin(), dy3.begin())); thrust::inclusive_scan(my_i_zip, my_i_zip+ds, my_o_zip, scan_op()); thrust::host_vector<mt> hx1 = dx1; thrust::copy_n(hx1.begin(), snip, std::ostream_iterator<mt>(std::cout, ",")); thrust::copy_n(hx1.begin()+ds-snip, snip, std::ostream_iterator<mt>(std::cout, ",")); std::cout << std::endl; } $ nvcc -std=c++14 t1930.cu -o t1930 $ cuda-memcheck ./t1930 ========= CUDA-MEMCHECK 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.218,601.275,576.315,607.993,582.947,614.621,589.516,621.699,595.644,628.843, 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.219,601.275,576.316,607.994,582.948,614.621,589.516,621.7,595.644,628.843, ========= ERROR SUMMARY: 0 errors $ 是的，上面有一些结果在第 6 位数字上有所不同。考虑到串行和并行方法之间非常不同的操作顺序，我将此归因于 float 分辨率的限制。如果您将 typedef 更改为 double，结果将看起来完全匹配。 既然您已经询问过它，这里有一个等效的实现，它是使用之前使用 cudaMalloc 分配的设备数据进行演示的： $ cat t1930.cu #include <iostream> #include <thrust/device_vector.h> #include <thrust/host_vector.h> #include <thrust/scan.h> #include <thrust/copy.h> #include <thrust/iterator/zip_iterator.h> #include <thrust/iterator/constant_iterator.h> #include <cstdlib> #include <cstdio> template <typename T> void cpufunction(T *result, T *oldArray, size_t size, T k){ for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } } struct scan_op // as per blelloch (1.7) { template <typename T1, typename T2> __host__ __device__ T1 operator()(const T1 &t1, const T2 &t2){ T1 ret; thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<2>(t2) + thrust::get<1>(t1)*thrust::get<4>(t2)+thrust::get<0>(t2); thrust::get<1>(ret) = thrust::get<0>(t1)*thrust::get<3>(t2) + thrust::get<1>(t1)*thrust::get<5>(t2)+thrust::get<1>(t2); thrust::get<2>(ret) = thrust::get<2>(t1)*thrust::get<2>(t2) + thrust::get<3>(t1)*thrust::get<4>(t2); thrust::get<3>(ret) = thrust::get<2>(t1)*thrust::get<3>(t2) + thrust::get<3>(t1)*thrust::get<5>(t2); thrust::get<4>(ret) = thrust::get<4>(t1)*thrust::get<2>(t2) + thrust::get<5>(t1)*thrust::get<4>(t2); thrust::get<5>(ret) = thrust::get<4>(t1)*thrust::get<3>(t2) + thrust::get<5>(t1)*thrust::get<5>(t2); return ret; } }; typedef double mt; const size_t ds = 512; const mt k = 1.01; const int snip = 10; int main(){ mt *b1 = new mt[ds]; // b as in blelloch (1.5) mt *cr = new mt[ds]; // cpu result for (int i = 0; i < ds; i++) { b1[i] = rand()/(float)RAND_MAX;} cr[0] = b1[0]; cr[1] = b1[1]; cpufunction(cr, b1, ds, k); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; mt *db; cudaMalloc(&db, ds*sizeof(db[0])); cudaMemcpy(db, b1, ds*sizeof(db[0]), cudaMemcpyHostToDevice); thrust::device_ptr<mt> dp_db = thrust::device_pointer_cast(db); auto b0 = thrust::constant_iterator<mt>(0); auto a0 = thrust::constant_iterator<mt>(0); auto a1 = thrust::constant_iterator<mt>(1); auto a2 = thrust::constant_iterator<mt>(k); auto a3 = thrust::constant_iterator<mt>(0); thrust::device_vector<mt> dx1(ds); thrust::device_vector<mt> dx0(ds); thrust::device_vector<mt> dy0(ds); thrust::device_vector<mt> dy1(ds); thrust::device_vector<mt> dy2(ds); thrust::device_vector<mt> dy3(ds); auto my_i_zip = thrust::make_zip_iterator(thrust::make_tuple(dp_db, b0, a0, a1, a2, a3)); auto my_o_zip = thrust::make_zip_iterator(thrust::make_tuple(dx1.begin(), dx0.begin(), dy0.begin(), dy1.begin(), dy2.begin(), dy3.begin())); thrust::inclusive_scan(my_i_zip, my_i_zip+ds, my_o_zip, scan_op()); cudaMemcpy(cr, thrust::raw_pointer_cast(dx1.data()), ds*sizeof(cr[0]), cudaMemcpyDeviceToHost); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; } $ nvcc -std=c++14 t1930.cu -o t1930 $ cuda-memcheck ./t1930 ========= CUDA-MEMCHECK 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.219,601.275,576.316,607.994,582.948,614.622,589.516,621.7,595.645,628.844, 0.840188,0.394383,1.63169,1.19677,2.55965,1.40629,2.92047,2.18858,3.22745,2.76443,570.219,601.275,576.316,607.994,582.948,614.622,589.516,621.7,595.645,628.844, ========= ERROR SUMMARY: 0 errors 这两种方法之间应该没有显着的性能差异。 （但是，在本例中，我碰巧将 typedef 切换为 double，因此这会有所不同。）使用 cudaMalloc 作为各种状态向量（device_vector、dx0）的 dx1 的替代品、 dy0、dy1 ...）可能会稍微快一些，因为 device_vector 首先执行 cudaMalloc 样式分配，然后启动内核将分配归零。对于状态向量来说，这个归零步骤是不必要的。如果您有兴趣，这里给出的模式应该演示如何做到这一点。 这是一个完全消除使用 thrust::device_vector 和 thrust::host_vector 的版本： #include <iostream> #include <thrust/device_ptr.h> #include <thrust/scan.h> #include <thrust/iterator/zip_iterator.h> #include <thrust/iterator/constant_iterator.h> #include <cstdlib> template <typename T> void cpufunction(T *result, T *oldArray, size_t size, T k){ for (int i = 2; i<size; i++) { result[i] = oldArray[i] + k * result[i-2]; } } struct scan_op // as per blelloch (1.7) { template <typename T1, typename T2> __host__ __device__ T1 operator()(const T1 &t1, const T2 &t2){ T1 ret; thrust::get<0>(ret) = thrust::get<0>(t1)*thrust::get<2>(t2) + thrust::get<1>(t1)*thrust::get<4>(t2)+thrust::get<0>(t2); thrust::get<1>(ret) = thrust::get<0>(t1)*thrust::get<3>(t2) + thrust::get<1>(t1)*thrust::get<5>(t2)+thrust::get<1>(t2); thrust::get<2>(ret) = thrust::get<2>(t1)*thrust::get<2>(t2) + thrust::get<3>(t1)*thrust::get<4>(t2); thrust::get<3>(ret) = thrust::get<2>(t1)*thrust::get<3>(t2) + thrust::get<3>(t1)*thrust::get<5>(t2); thrust::get<4>(ret) = thrust::get<4>(t1)*thrust::get<2>(t2) + thrust::get<5>(t1)*thrust::get<4>(t2); thrust::get<5>(ret) = thrust::get<4>(t1)*thrust::get<3>(t2) + thrust::get<5>(t1)*thrust::get<5>(t2); return ret; } }; typedef float mt; const size_t ds = 32768*4; const mt k = 1.001; const int snip = 10; int main(){ mt *b1 = new mt[ds]; // b as in blelloch (1.5) mt *cr = new mt[ds]; // result for (int i = 0; i < ds; i++) { b1[i] = (rand()/(float)RAND_MAX)-0.5;} cr[0] = b1[0]; cr[1] = b1[1]; cpufunction(cr, b1, ds, k); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; mt *db, *dstate; cudaMalloc(&db, ds*sizeof(db[0])); cudaMalloc(&dstate, 6*ds*sizeof(dstate[0])); cudaMemcpy(db, b1, ds*sizeof(db[0]), cudaMemcpyHostToDevice); thrust::device_ptr<mt> dp_db = thrust::device_pointer_cast(db); auto b0 = thrust::constant_iterator<mt>(0); auto a0 = thrust::constant_iterator<mt>(0); auto a1 = thrust::constant_iterator<mt>(1); auto a2 = thrust::constant_iterator<mt>(k); auto a3 = thrust::constant_iterator<mt>(0); thrust::device_ptr<mt> dx1 = thrust::device_pointer_cast(dstate); thrust::device_ptr<mt> dx0 = thrust::device_pointer_cast(dstate+ds); thrust::device_ptr<mt> dy0 = thrust::device_pointer_cast(dstate+2*ds); thrust::device_ptr<mt> dy1 = thrust::device_pointer_cast(dstate+3*ds); thrust::device_ptr<mt> dy2 = thrust::device_pointer_cast(dstate+4*ds); thrust::device_ptr<mt> dy3 = thrust::device_pointer_cast(dstate+5*ds); auto my_i_zip = thrust::make_zip_iterator(thrust::make_tuple(dp_db, b0, a0, a1, a2, a3)); auto my_o_zip = thrust::make_zip_iterator(thrust::make_tuple(dx1, dx0, dy0, dy1, dy2, dy3)); thrust::inclusive_scan(my_i_zip, my_i_zip+ds, my_o_zip, scan_op()); cudaMemcpy(cr, dstate, ds*sizeof(cr[0]), cudaMemcpyDeviceToHost); for (int i = 0; i < snip; i++) std::cout << cr[i] << ","; for (int i = ds-snip; i < ds; i++) std::cout << cr[i] << ","; std::cout << std::endl; } 这里是一些 CPU 代码，显示了源自 https://www.cs.cmu.edu/~guyb/papers/Ble93.pdf 的公式的可能实现，以将高阶递归表示为扫描操作。 关键思想是扫描结果的每个元素不是标量，而是包含前 n 个标量结果的向量。这样，扫描运算符中就可以使用所有必需的先前结果来计算下一个结果。 #include <iostream> #include <algorithm> #include <numeric> #include <array> void calculate1(std::vector<int> vec, int k){ std::vector<int> result(vec.size(), 0); for(int i = 2; i < vec.size(); i++){ result[i] = vec[i] + k * result[i-2]; } std::cerr << "calculate1 result: "; for(auto x : result){ std::cerr << x << ", "; } std::cerr << "\n"; } struct S{ //data[0] stores result of last iteration //data[1] stores result of second last iteration std::array<int, 2> data; }; std::ostream& operator<<(std::ostream& os, S s){ os << "(" << s.data[0] << "," << s.data[1] << ")"; } void calculate2(std::vector<int> vec, int k){ S initvalue{{0,0}}; std::vector<S> result(vec.size(), initvalue); std::exclusive_scan( vec.begin() + 2, vec.end(), result.begin(), initvalue, [k](S left, int right){ S result; /*A = ( 0 1 k 0 ) Compute result = left * A + (right 0) */ result.data[0] = right + k * left.data[1]; result.data[1] = left.data[0]; return result; } ); std::cerr << "calculate2 result: "; for(auto x : result){ std::cerr << x << ", "; } std::cerr << "\n"; } int main(){ const int k = 5; const std::vector<int> vec1{1,3,5,7,9,11,3,6,7,1,2,4}; calculate1(vec1, k); calculate2(vec1, k); } https://godbolt.org/z/cszzn8Ec8 输出： calculate1 result: 0, 0, 5, 7, 34, 46, 173, 236, 872, 1181, 4362, 5909, calculate2 result: (0,0), (5,0), (7,5), (34,7), (46,34), (173,46), (236,173), (872,236), (1181,872), (4362,1181), (0,0), (0,0), 某处仍然存在逐一错误，但人们可以理解其背后的想法。 我之前说过这种方法可以用于 CUDA 中的并行扫描。这是不正确的。对于并行扫描，扫描算子必须具有一个附加属性，即结合性，即 (a OP b) OP c == a OP (b OP c)。这种方法的情况并非如此。 Robert Crovella 的答案展示了如何导出可用于并行扫描的关联扫描算子。

thrust::transform() 导致从主机到设备的 cudaErrorIllegalAddress

以下test.cu程序 #包括 #包括 #包括 #包括 #包括 以下test.cu程序 #include <thrust/copy.h> #include <thrust/execution_policy.h> #include <thrust/transform.h> #include <thrust/device_vector.h> #include <thrust/host_vector.h> #include <iostream> using HOST_TYPE=int32_t; using DEVICE_TYPE=int; template <typename T> struct Cast { __host__ __device__ T operator()(HOST_TYPE i) const { return static_cast<T>(i); } }; int main() { // Initialize host data thrust::host_vector<HOST_TYPE> const h_vec{1, 2, 3, 4, 5}; // Allocate space on the device thrust::device_vector<DEVICE_TYPE> device_data(h_vec.size()); // Copy data from host to device //thrust::copy(h_vec.cbegin(), h_vec.cend(), device_data.begin()); // this works thrust::transform(h_vec.cbegin(), h_vec.cend(), device_data.begin(), Cast<DEVICE_TYPE>{}); // Copy back to host to check thrust::host_vector<DEVICE_TYPE> host_data_copy = device_data; for (DEVICE_TYPE val : host_data_copy) { std::cout << val << " "; } std::cout << std::endl; return 0; } 原因 $ nvcc test.cu $ ./a.out terminate called after throwing an instance of 'thrust::system::system_error' what(): parallel_for: failed to synchronize: cudaErrorIllegalAddress: an illegal memory access was encountered Aborted (core dumped) 这发生在生产线上 thrust::transform(h_vec.cbegin(), h_vec.cend(), device_data.begin(), Cast<DEVICE_TYPE>{}); 即使类似的thrust::copy()运行良好： thrust::copy(h_vec.cbegin(), h_vec.cend(), device_data.begin()); // this works 我在文档中找不到任何内容说明thrust::transform()不应在设备和主机之间转换数据。我是不是在某个地方错过了这个？ 使用 thrust::host 或 thrust::device 执行策略没有帮助。 版本： $ nvcc --version nvcc: NVIDIA (R) Cuda compiler driver Copyright (c) 2005-2023 NVIDIA Corporation Built on Fri_Nov__3_17:16:49_PDT_2023 Cuda compilation tools, release 12.3, V12.3.103 Build cuda_12.3.r12.3/compiler.33492891_0 注意：实际应用需要 HOST_TYPE=char，但为了调试/说明目的，将其更改为 HOST_TYPE=int32_t，并与 std::copy() 进行比较。 请参阅转换： 除了可以在主机和设备之间复制数据的 Thrust::copy 之外，Thrust 算法的所有迭代器参数都应该位于同一位置：要么全部位于主机上，要么全部位于设备上。当违反此要求时，编译器将产生错误消息。

获取 CUDA Thrust 以使用您选择的 CUDA 流

查看 CUDA Thrust 代码中的内核启动，似乎它们总是使用默认流。我可以让 Thrust 使用我选择的流吗？我是否遗漏了 API 中的某些内容？