OpenCL 和 CUDA 中的持久线程

我读过一些关于 GPGPU 的“持久线程”的论文，但我不太理解它。任何人都可以给我一个例子或向我展示这种编程方式的使用吗？

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

持久线程只不过是一个 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

CUDA 利用单指令多数据 (SIMD) 编程模型。计算线程以块的形式组织，并且线程块被分配给不同的流多处理器（SM）。 SM 上线程块的执行是通过将线程排列在warps of 32线程：每个线程束以锁步方式运行并准确执行 对不同数据的相同指令。

一般来说，为了填满 GPU，内核会启动更多的块，这些块实际上可以托管在 SM 上。由于并非所有块都可以托管在 SM 上，因此工作调度程序会在块完成计算时执行上下文切换。应该注意的是，块的切换完全由调度程序在硬件中管理，程序员无法影响块如何调度到 SM 上。这暴露了所有那些不完全适合 SIMD 编程模型并且存在工作不平衡的算法的局限性。确实是一个块A不会被另一个块取代B在同一个 SM 上，直到块的最后一个线程A将尚未完成执行。

尽管 CUDA 没有向程序员公开硬件调度程序，持久线程style 通过依赖工作队列绕过硬件调度程序。当一个块完成时，它会检查队列中是否有更多工作，并继续这样做，直到没有剩余工作，此时该块将退出。通过这种方式，内核启动时会使用与可用 SM 数量相同的块。

The 持久线程下面的示例可以更好地说明该技术，该示例取自演示文稿

“GPGPU”计算和CUDA/OpenCL编程模型

论文中提供了另一个更详细的示例

了解 GPU 上光线遍历的效率

// 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);