我正在编写一个应用程序,它使用多个并发CUDA流。当我的thrust::reduce_by_key
调用似乎要写入可分页内存时,我的其他流正在阻塞。我认为返回值是问题所在。
如何防止返回值被写入可分页内存?
我将包括演示我尝试的解决方案的代码。
#include <thrust/system/cuda/vector.h>
#include <thrust/host_vector.h>
#include <thrust/pair.h>
#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>
#include <thrust/random.h>
int main(void)
{
int N = 20;
thrust::default_random_engine rng;
thrust::uniform_int_distribution<int> dist(10, 99);
// initialize data
thrust::device_vector<int> array(N);
for (size_t i = 0; i < array.size(); i++)
array[i] = dist(rng);
// allocate storage for sums and indices
thrust::device_vector<int> sums(N);
thrust::device_vector<int> indices(N);
// make a pinned memory location for the returned pair of iterators
typedef thrust::device_vector<int>::iterator dIter;
thrust::pair<dIter, dIter>* new_end;
const unsigned int bytes = sizeof(thrust::pair<dIter, dIter>);
cudaMallocHost((void**)&new_end, bytes);
for(int i = 0 ; i< 20; i++){ // you can see in the profiler each operator writes 4 bytes to pageable memory
*new_end = thrust::reduce_by_key
(thrust::make_counting_iterator(0),
thrust::make_counting_iterator(N),
array.begin(),
indices.begin(),
sums.begin(),
thrust::equal_to<int>(),
thrust::plus<int>());
}
std::cout << "done \n";
return 0;
}
这是我的分析器的图片,显示了从设备到主机可分页内存
的拷贝
1条答案
按热度按时间t40tm48m1#
我正在编写一个应用程序,它使用多个并发CUDA流。当我的
thrust::reduce_by_key
似乎要写入可分页内存时,我的其他流正在阻塞这种阻塞行为不是由“写入可分页内存”引起的。它是由显式同步调用引起。一般来说,自CUDA 10.1(Thrust 1.9.4)版本起,all normal synchronous algorithms are blocking。您可以通过使用探查器检查API跟踪来自己确认这一点。但是,您至少可以通过将调用启动到stream中来限制阻塞的范围,尽管我懒得测试这是否以有用的方式修改了
cuda_cub::synchronize
的行为。如何防止返回值被写入可分页内存?
不是说这和你的问题有任何关系,但你不能。重要的是要记住,与您最初的问题相反,
thrust::reduce_by_key
不是内核,它是执行一系列操作的主机代码,包括将返回值从设备内存复制到主机堆栈变量。程序员无法控制内部机制,显然,试图使用自己的固定内存值来接受通过值传递的结果是荒谬的,不会有任何效果。正如评论中所建议的,如果您需要您的问题所建议的操作内部控制的粒度级别,那么推力是错误的选择。使用cub::device::reduce_by_key --这与thrust实现使用的算法相同,但您可以显式控制临时内存、同步、流以及如何访问调用结果。然而,这并不适合初学者。