当我仍然有足够的内存时，cuda push::sort 遇到了内存问题

Posted 2023-03-23

【中文标题】当我仍然有足够的内存时，cuda push::sort 遇到了内存问题

【英文标题】：cuda thrust::sort met memory problem when I still have enough memory

【发布时间】：2021-02-03 02:02:06

【问题描述】：

我在 ubuntu18.04 上使用 cuda10.2。我的gpu是tesla T4，内存16G，目前gpu上没有运行其他程序。 一小段代码如下：

#include <iostream>
#include <algorithm>
#include <random>
#include <vector>
#include <numeric>
#include <algorithm>
#include <chrono>

#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <thrust/execution_policy.h>


struct sort_functor 

    thrust::device_ptr<float> data;
    int stride = 1;
    __host__ __device__
    void operator()(int idx) 
        thrust::sort(thrust::device,
                data + idx * stride, 
                data + (idx + 1) * stride);
    
;


int main() 
    std::random_device rd;
    std::mt19937 engine;
    engine.seed(rd());
    std::uniform_real_distribution<float> u(0, 90.);

    int M = 8;
    int N = 8 * 384 * 300;

    std::vector<float> v(M * N);
    std::generate(v.begin(), v.end(), [&]()return u(engine););
    thrust::host_vector<float> hv(v.begin(), v.end());
    thrust::device_vector<float> dv = hv;

    thrust::device_vector<float> res(dv.begin(), dv.end());

    thrust::device_vector<int> index(M);
    thrust::sequence(thrust::device, index.begin(), index.end(), 0, 1);

    thrust::for_each(thrust::device, index.begin(), index.end(), 
            sort_functorres.data(), N
            );
    cudaDeviceSynchronize();

    return 0;

错误信息是：

temporary_buffer::allocate: get_temporary_buffer failed
temporary_buffer::allocate: get_temporary_buffer failed
temporary_buffer::allocate: get_temporary_buffer failed
temporary_buffer::allocate: get_temporary_buffer failed
temporary_buffer::allocate: get_temporary_buffer failed
temporary_buffer::allocate: get_temporary_buffer failed
terminate called after throwing an instance of 'thrust::system::system_error'
  what():  for_each: failed to synchronize: cudaErrorLaunchFailure: unspecified launch failure
Aborted (core dumped)

请问我该如何解决这个问题？

【问题讨论】：

不要对设备向量使用执行策略

@talonmies 谢谢你告诉我这个，你能写一个答案并解释原因吗？

这不是解决您问题的方法，而是一种观察。 Thrust 自动知道如何处理主机代码中主机和设备向量的执行策略。唯一需要执行指针的情况是将指针（没有基于标签的元数据）传递给推力调用。然后你需要编译器的执行策略知道如何将调用分派到正确的后端

【参考方案1】：

thrust::sortrequires O(N) temporary memory allocation。当您从设备代码（在您的仿函数中）调用它时，临时内存分配（对于每个调用 - 即从您的 8 个调用中的每一个调用）将在设备上使用 new 或 malloc 在引擎盖下完成，并且分配将来自“设备堆”空间。 device heap space is by default limited to 8MB，但您可以更改它。您已达到此限制。

如果您在 main 例程的顶部添加以下内容：

cudaError_t err = cudaDeviceSetLimit(cudaLimitMallocHeapSize, 1048576ULL*1024);

您的代码运行时对我来说没有任何运行时错误。

我并不是说我仔细计算了上面的 1GB 值。我只是选择了一个比 8MB 大得多但比 16GB 小得多的值，它似乎可以工作。在一般情况下，您可能应该carefully estimate 您需要的临时分配大小。