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 中线程 channel 号的位置。您发布的内联 PTX 只是读取该寄存器的值并将其存储在一个变量中，您可以在该变量中使用您发布的后续 C++ 代码。

每个版本的 CUDA 工具包都附带一个 PTX 汇编语言引用指南，您可以使用它来查找此类内容。