CUDA5.0样本AdvancedQuickSort

Question

我正在阅读CUDA 5.0样本(AdvancedQuickSort).但是,由于以下代码,我完全无法理解此示例:

// Now compute my own personal offset within this. I need to know how many
// threads with a lane ID less than mine are going to write to the same buffer
// as me. We can use popc to implement a single-operation warp scan in this case.
unsigned lane_mask_lt;
asm( "mov.u32 %0, %%lanemask_lt;" : "=r"(lane_mask_lt) );
unsigned int my_mask = greater ? gt_mask : lt_mask;
unsigned int my_offset = __popc(my_mask & lane_mask_lt);

这是__global__ void qsort_warp函数中的函数,尤其是代码中的汇编语言.任何人都可以帮我解释这种汇编语言的含义吗？

Answer 1

%lanemask_lt是PTX程序集中的一个特殊的只读寄存器,它使用32位掩码进行初始化,其位数设置的位置小于warp中线程的通道号.您发布的内联PTX只是读取该寄存器的值并将其存储在变量中,可以在您发布的后续C++代码中使用该变量.

每个版本的CUDA工具包都附带一个PTX程序集语言参考指南,您可以使用它来查找这样的内容.