OpenCL和CUDA中的持久线程

Question

我读了一些关于GPGPU的“持久线程（persistent threads）”的论文，但我并不真正理解它。有人能给我举个例子或展示这种编程方式的使用吗？

在阅读和搜索“持久线程”后，我记住的是：

持久线程不过是一个while循环，使线程保持运行并计算许多工作。

这正确吗？谢谢。

参考：http://www.idav.ucdavis.edu/publications/print_pub?pub_id=1089 http://developer.download.nvidia.com/GTC/PDF/GTC2012/PresentationPDF/S0157-GTC2012-Persistent-Threads-Computing.pdf

- AmineMs

---

也许你应该包含一些具体的论文参考，谈论一下“持久线程”。 - Robert Crovella

持久线程风格GPU编程在GPGPU工作负载中的研究 - AmineMs

如果您观看演示文稿，您可能会获得一些额外的有用见解。我的理解是，它基本上就像您所描述的那样，意味着一个不退出并持续从全局内存队列中轮询工作的内核。 - Robert Crovella

Answer 1

CUDA采用单指令多数据（SIMD）编程模型。计算线程以块的形式组织，并将线程块分配给不同的流式多处理器（SM）。在线程块上执行SM时，通过将线程分组为线程束（warps），每个线程束中有32个线程：每个线程束都是锁步运行的，并在不同的数据上执行完全相同的指令。

通常，为了填满GPU，核函数会启动比实际可承载在SM上的块数多得多的块。由于并非所有块都可以托管在SM上，因此工作调度程序在块完成计算时执行上下文切换。应该注意的是，块的切换完全由调度程序的硬件管理，并且程序员无法影响块如何被调度到SM上。这暴露了对于那些不完全适合SIMD编程模型且存在工作负载不平衡的所有算法的限制。事实上，在块A的最后一个线程没有执行完之前，块A不会被同一SM上的另一个块B替换。

虽然CUDA没有向程序员公开硬件调度程序，但持久线程技术通过依赖工作队列来绕过硬件调度程序。当一个块完成时，它检查队列是否还有更多工作要做，并继续这样做直到没有工作为止，然后该块退役。通过这种方式，核函数启动的块数与可用SM数量相同。

持久线程技术可以通过以下示例更好地说明，该示例摘自演示文稿：

“GPGPU” computing and the CUDA/OpenCL Programming Model

另一个更详细的示例在以下论文中提供：

Understanding the efficiency of ray traversal on GPUs

// Persistent thread: Run until work is done, processing multiple work per thread
// rather than just one. Terminates when no more work is available

// count represents the number of data to be processed

__global__  void persistent(int* ahead, int* bhead, int count, float* a, float* b)
{
    int local_input_data_index, local_output_data_index;
while ((local_input_data_index = read_and_increment(ahead)) <   count)
{                                   
        load_locally(a[local_input_data_index]);

        do_work_with_locally_loaded_data();

        int out_index = read_and_increment(bhead);

        write_result(b[out_index]);
    }
}

// Launch exactly enough threads to fill up machine (to achieve sufficient parallelism 
// and latency hiding)
persistent<<numBlocks,blockSize>>(ahead_addr, bhead_addr, total_count, A, B);

- Vitality

---

投票赞同以强调确保启动足够的线程来充分利用硬件资源。 - biubiuty

Answer 2

非常易于理解。通常每个工作项处理少量工作。如果想节省工作组切换时间，可以让一个工作项使用循环处理大量工作。例如，您有一张图片，它是1920x1080，您有1920个工作项，每个工作项使用循环处理1080像素的一列。

- Hunter Wang