cuda - CUDA コードのプロファイリング: 合体メモリ読み取りでの予期しない命令カウント

Question

小さな入力データ (= 512 要素) の非常にダンプされた並べ替えアルゴリズムをプロファイリングしています。合体した構造体の配列を読み取るカーネルを呼び出しています。

構造体は次のようになります。

struct __align__(8) Elements 
{
     float weight;
     int value;
};

nvprof は、L1 ミス/ヒットおよび gdl 命令について次の命令カウントを提供します。

                  Invocations    Avg       Min       Max  Event Name
        Kernel: sort(Elements*)
                      500         0         0         0  gld_inst_8bit
                      500         0         0         0  gld_inst_16bit
                      500      1024      1024      1024  gld_inst_32bit
                      500         0         0         0  gld_inst_64bit
                      500         0         0         0  gld_inst_128bit
                      500       120       120       120  l1_global_load_hit
                      500       120       120       120  l1_global_load_miss
                      500         0         0         0  uncached_global_load_tr.

次のように構造体のレイアウトを変更すると:

struct __align__(8) Elements 
{
     float weight;
     float value;
};

プロファイリングの出力は次のようになります。

                  Invocations    Avg       Min       Max  Event Name
Device 0
        Kernel: sort(Elements*)
                      500         0         0         0  gld_inst_8bit
                      500         0         0         0  gld_inst_16bit
                      500         0         0         0  gld_inst_32bit
                      500       512       512       512  gld_inst_64bit
                      500         0         0         0  gld_inst_128bit
                      500         0         0         0  l1_global_load_hit
                      500       120       120       120  l1_global_load_miss
                      500         0         0         0  uncached_global_load_tr.

実行時間への影響はまったくありませんが、GPU がコードの最初のバリアントで 32 ビットのロード命令を実行し、2 番目のバリアントで 64 ビットの命令を実行する理由がわかりません。

カーネルは 1 ブロックと 512 スレッドで呼び出されます (そのため、l1_global_load_x カウンターが正しくない可能性があります)。すべてが CUDA 5.0 を搭載した GeForce 480 で行われます。

EDIT： ソートカーネル（少し短縮）：

__global__ void sort(Elements* nearest)
{
    ThreadIndex idx = index();

    __shared__ Elements temp[MAX_ELEMENTS];
    __shared__ int index_cache[MAX_ELEMENTS];

    temp[idx.x] = nearest[idx.x];

    WeightedElements elem = temp[idx.x];
    __syncthreads();

    int c = 0;

    // some index crunching 

    nearest[idx.x] = temp[c];
}