9

我已经阅读了一些关于 GPGPU 的“持久线程”的论文，但我并不真正理解它。谁能给我一个例子或向我展示这种编程方式的用法？

在阅读和谷歌搜索“持久线程”后，我记住了什么：

Presistent 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

4

2 回答 2

11

CUDA 利用单指令多数据 (SIMD) 编程模型。计算线程按块组织，线程块分配给不同的流式多处理器 (SM)。SM 上线程块的执行是通过将线程安排在线程束中来执行的32：每个束以锁步方式运行，并对不同的数据执行完全相同的指令。

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

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

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

“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);
于 2014-06-10T17:06:53.367 回答
4

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

于 2014-01-20T07:45:45.897 回答