sort_by_key in thrust 是阻塞调用吗？

Question

我反复排队一系列内核：

for 1..100:
    for 1..10000:
        // Enqueue GPU kernels
        Kernel 1 - update each element of array  
        Kernel 2 - sort array  
        Kernel 3 - operate on array  
    end
    // run some CPU code
    output "Waiting for GPU to finish"
    // copy from device to host
    cudaMemcpy ... D2H(array)
end

内核 3 的阶数为 O(N^2)，因此是迄今为止最慢的。对于内核 2，我直接在设备上使用 thrust::sort_by_key：

thrust::device_ptr<unsigned int> key(dKey);
thrust::device_ptr<unsigned int> value(dValue);
thrust::sort_by_key(key,key+N,value);

似乎这个对 thrust 的调用是阻塞的，因为 CPU 代码只有在内部循环完成后才会执行。我看到这个是因为如果我删除对 sort_by_key 的调用，主机代码（正确地）在内部循环完成之前输出 "Waiting" 字符串，而如果我运行排序则不会。

有没有办法异步调用thrust::sort_by_key？

Answer 1

首先考虑有一个内核启动队列，它只能容纳这么多等待启动。一旦启动队列已满，任何类型 的额外内核启动都会被阻塞。在空队列槽可用之前，主机线程不会继续（超出那些启动请求）。我很确定 3 个内核启动的 10000 次迭代将在达到 10000 次迭代之前填满此队列。因此，如果您按顺序启动 30000 个内核，那么任何类型的重要内核启动都会有一些延迟（我认为）。（然而，最终，当所有内核都添加到队列中时，因为有些内核已经完成，如果没有其他阻塞行为，那么在所有内核实际完成之前，您会看到 "waiting..." 消息。）
thrust::sort_by_key requires temporary storage（大小约等于您的数据集大小）。每次使用时，都会通过 cudaMalloc 操作在后台分配此临时存储空间。这个cudaMalloc操作是阻塞。当从主机线程启动 cudaMalloc 时，它会在内核 activity 中等待间隙，然后才能继续。

要解决第 2 项，似乎至少有 2 种可能的方法：

提供一个thrust custom allocator。根据此分配器的特性，您可能能够消除阻塞 cudaMalloc 行为。（但请参阅下面的讨论）
使用cub SortPairs。这里的优点（如我所见 - 你的例子不完整）是你可以进行一次分配（假设你知道整个循环迭代中最坏情况下的临时存储大小）并且无需在你的内部进行临时内存分配环形。

据我所知，thrust 方法（上面的 1）仍然会在每次迭代时有效地执行某种临时 allocation/free 步骤，即使您提供自定义分配器也是如此。如果您有一个设计良好的自定义分配器，那么这可能几乎是一个 "no-op"。 cub 方法似乎有需要知道最大大小的缺点（为了完全消除对 allocation/free 步骤的需要），但我认为相同的要求将适用于推力自定义分配器。否则，如果您需要在某个时候分配更多内存，自定义分配器实际上将不得不执行类似 cudaMalloc 的操作，这会在工作中造成麻烦。

sort_by_key in thrust 是阻塞调用吗？

Is sort_by_key in thrust a blocking call?

cuda

gpgpu

thrust