在OpenCL和CUDA线程持续(Persistent threads in OpenCL and

我看过一些论文谈到“老大难线程”的GPGPU，但我真的不明白。任何一个可以给我一个例子或告诉我使用这种编程方式的？

我一直在我脑海阅读和谷歌上搜索“持久的主题”后：

Presistent的思路，没有比一个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

Answer 1:

CUDA利用了单指令多数据（SIMD）编程模型。计算线程块组织和线程块被分配到不同的流式多处理器（SM）。上的SM线程块的执行是通过在经纱的排列线程执行32线程：每根经在锁步操作，并执行准确的不同数据的相同指令。

一般来说，填补了GPU，内核与实际能够在SMS托管更多块启动。由于不是所有的块可以在SM被托管，作业调度器执行当一个块已完成计算上下文切换。值得注意的是，块的切换是完全在硬件由调度管理，以及程序员影响块被调度到SM的手段。这暴露了所有那些不完全符合一个SIMD编程模型算法的限制，且有就是工作不平衡。事实上，一个块A不会被另一块代替B在同一个SM，直到块的最后一个线程A不会完成执行。

虽然CUDA不公开硬件调度给程序员， 持久线程风格依靠工作队列绕过硬件调度器。当一个块结束时，会检查更多的工作队列，并继续这样做，直到没有工作留下，此时块退休。通过这种方式，内核与尽可能多块可用短信号码发布。

持久线程技术更好通过下列实施例，其已经从呈现采取所示

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

Answer 2:

很容易理解。一般每个工作项处理工作的量小。如果你要救救工作组交换机的时候，你可以让一个工作项目过程中使用循环了很多工作。例如，你有一个形象，这是1920×1080，你有1920的工作项，每个工作项的处理循环使用1080像素的一列。