cuda - 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は、ワープ内のスレッドのレーン番号よりも小さい位置にビットが設定された32ビットマスクで初期化される、PTXアセンブリの特別な読み取り専用レジスタです。投稿したインラインPTXは、そのレジスタの値を読み取り、それを変数に格納して、投稿した後続のC++コードで使用できるようにするだけです。

CUDAツールキットのすべてのバージョンには、このようなものを検索するために使用できるPTXアセンブリ言語リファレンスガイドが付属しています。