OpenCL和CUDA中的持久线程

Question

我已经阅读了一些关于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

Answer 1

CUDA利用单指令多数据(SIMD)编程模型.计算线程以块的形式组织,并且线程块被分配给不同的流式多处理器(SM).通过在线程的经线中布置线程来执行SM上的线程块的执行32:每个warp以锁定步骤操作并且对不同数据执行完全相同的指令.

通常,为了填充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);