第一次从推力执行排序需要太长时间

Question

我正在使用 opengl（实现 sph 算法）开发流体模拟器。我已经尝试了很多方法来运行我的模拟器，首先我使用八叉树，然后是哈希图，现在我正在尝试使用 Z 顺序，为此我需要根据它们的索引对我的粒子进行排序。

我无法理解的事实是，如果我有一个 thrust::sort 需要 15 毫秒，如果我有两个 thrust::sort 则需要 17 毫秒。

为了进一步说明，我在 opengl 中做我的模拟器（我所有的缓冲区都是使用 opengl 创建的），我使用 cuda 互操作来使用 cuda 对我的缓冲区进行排序。

这是我获取缓冲区并将它们“link”到 cuda

的部分

//I use this if to do the registerBuffer only one time
if (first == 0) {
        //index
        IBuffer* bIndex = RESOURCEMANAGER->getBuffer("particleLib::Index");
        int buffIdIndex = bIndex->getPropi(IBuffer::ID);
        //Position
        IBuffer* bPosition = RESOURCEMANAGER->getBuffer("particleLib::Position");
        int buffIdPosition = bPosition->getPropi(IBuffer::ID);
        //TempIndex
        IBuffer* bTempIndex = RESOURCEMANAGER->getBuffer("particleLib::TempIndex");
        int buffIdTempIndex = bTempIndex->getPropi(IBuffer::ID);
        //Velocity
        IBuffer* bVelocity = RESOURCEMANAGER->getBuffer("particleLib::Velocity");
        int buffIdVelocity = bVelocity->getPropi(IBuffer::ID);

        // register this buffer object with CUDA
        //So devia chamar isto uma vez
        cudaGraphicsGLRegisterBuffer(&cuda_ssbo_Index, buffIdIndex, cudaGraphicsMapFlagsNone);
        cudaGraphicsGLRegisterBuffer(&cuda_ssbo_TempIndex, buffIdTempIndex, cudaGraphicsMapFlagsNone);
        cudaGraphicsGLRegisterBuffer(&cuda_ssbo_Position, buffIdPosition, cudaGraphicsMapFlagsNone);
        cudaGraphicsGLRegisterBuffer(&cuda_ssbo_Velocity, buffIdVelocity, cudaGraphicsMapFlagsNone);
        first = 1;
    }
    

    // map OpenGL buffer object for writing from CUDA
    int* dptrssboIndex;
    int* dptrssboTempIndex;
    float4 * dptrssboPosition;
    float4 * dptrssboVelocity;

    cudaGraphicsMapResources(1, &cuda_ssbo_Index, 0);
    cudaGraphicsMapResources(1, &cuda_ssbo_TempIndex, 0);
    cudaGraphicsMapResources(1, &cuda_ssbo_Position, 0);
    cudaGraphicsMapResources(1, &cuda_ssbo_Velocity, 0);

    size_t num_bytesssbo_Index;
    size_t num_bytesssbo_TempIndex;
    size_t num_bytesssbo_Position;
    size_t num_bytesssbo_Velocity;

    cudaGraphicsResourceGetMappedPointer((void**)&dptrssboIndex, &num_bytesssbo_Index, cuda_ssbo_Index);
    cudaGraphicsResourceGetMappedPointer((void**)&dptrssboTempIndex, &num_bytesssbo_TempIndex, cuda_ssbo_TempIndex);
    cudaGraphicsResourceGetMappedPointer((void**)&dptrssboPosition, &num_bytesssbo_Position, cuda_ssbo_Position);
    cudaGraphicsResourceGetMappedPointer((void**)&dptrssboVelocity, &num_bytesssbo_Velocity, cuda_ssbo_Velocity);

    mysort(&dptrssboIndex,&dptrssboPosition, &dptrssboTempIndex, &dptrssboVelocity,216000);

    cudaGraphicsUnmapResources(1, &cuda_ssbo_Index, 0);
    cudaGraphicsUnmapResources(1, &cuda_ssbo_TempIndex, 0);
    cudaGraphicsUnmapResources(1, &cuda_ssbo_Position, 0);
    cudaGraphicsUnmapResources(1, &cuda_ssbo_Velocity, 0);

这是来自 mysort 的代码

void mysort(int ** index1, float4 ** values1, int** index2, float4 ** values2,int particles){
    
    thrust::device_ptr<int> i1buff = thrust::device_pointer_cast(*(index1));
    thrust::device_ptr<float4> v1buff = thrust::device_pointer_cast(*(values1));
    thrust::device_ptr<int> i2buff = thrust::device_pointer_cast(*(index2));
    thrust::device_ptr<float4> v2buff = thrust::device_pointer_cast(*(values2));

    //sorts
    thrust::stable_sort_by_key(i1buff, i1buff + particles,v1buff); // 15 ms
    //cudaThreadSynchronize();
    thrust::stable_sort_by_key(i2buff, i2buff + particles, v2buff); // 17 ms


    //repetido so para ver o tempo
    thrust::stable_sort_by_key(i1buff, i1buff + particles, v1buff);
    //cudaThreadSynchronize();
    thrust::stable_sort_by_key(i2buff, i2buff + particles, v2buff); //4 sorts -> 19 ms

    //cudaThreadSynchronize();
}

谁能解释一下这是怎么回事？

编辑1：我使用 cudaDeviceSynchronize() 来测量每次排序所需的时间（如@Jérôme-Richard 所示），并且即使我更改订单，第一次排序总是需要更长的时间。另一个事实是，如果我的相机离场景更近，第一种需要更长的时间，这表明 Cuda 可能正在等待 opengl 完成他的工作，从而使第一种“需要更长的时间”。我也尝试对我的 mysort() 函数不进行排序，我唯一拥有的是 cudaDeviceSynchronize() 并且它花费了 15 毫秒，这再次表明 cuda 可能正在等待 opengl 完成最后一帧的工作.

编辑2：我做了更多的调试，我认为似乎是真的。真正的减速来自 cudaGraphicsMapResources 调用。据此(cudaGraphicsMapResources slow speed when mapping DirectX texture):

This function provides the synchronization guarantee that any graphics calls issued before cudaGraphicsMapResources() will complete before any subsequent CUDA work issued in stream begins.

是的，它正在等待 opengl 绘制一些东西，因为相机距离会影响 cudaGraphicsMapResources 花费的时间。

Answer 1

两个要点可以解释你的观察：

第一个 CUDA 函数调用隐式初始化运行时（相当慢）。
要排序的数组的实际内容can/often 影响排序的性能（关于 Thrust 实现中使用的算法）。数据排序后，可以更快地排序，因为它们已经排序。
Thrust 在许多提供的函数中进行少量同步（即它调用 cudaDeviceSynchronize），以确保可以从 CPU 端安全地读取从 GPU 传输的返回数据。当提交关于计算数据结果的多个相互依赖的 CUDA 内核时，它也在内部使用这种同步（您可以使用 Nvidia 分析器看到这一点）。关于在此函数之前进行的先前异步 CUDA 调用，过度同步会增加不必要的开销。

第一次从推力执行排序需要太长时间

First time executing sort from thrust it takes too long

c++

opengl

cuda

thrust