2

データ型のサイズが変わったときにカーネルで実行される命令の数を調べてみました

カスタムサイズのデータ​​構造を取得するために、次のように構造体を作成しました。

#define DATABYTES 40

__host__ __device__
struct floatArray
{
    float a[DATABYTES/4];
};

そして、上記のデータ型配列をある配列から別の配列にコピーするためだけにカーネルを作成しました

__global__
void copy_large_data(floatArray * d_in, floatArray * d_out)
{
    d_out[threadIdx.x] = d_in[threadIdx.x];
}

次に、単一のブロックで 32 スレッドのみに対して上記のカーネルを呼び出しました

floatArray * d_in;
floatArray * d_out;

cudaMalloc(&d_in, 32 * sizeof(floatArray));
cudaMalloc(&d_out, 32 * sizeof(floatArray));

copy_large_data<<<1, 32>>>(d_in, d_out);

を使用してプログラムをプロファイリングし、nvprofをチェックするinstructions per warpと、 の値の変化に伴ってパラメータ値が変化することがわかりましたDATABYTES。

私の質問は、この命令数の増加の理由がfloatArray構造体内の配列によるものかどうかです。aカーネルでコピーを呼び出すと、構造体内の配列の各要素が実際に展開およびコピーされfloatArray、より多くの命令が作成されるためです。

単一の命令を使用してカーネル内のカスタム構造体変数をコピーする方法はありますか?

4

1 に答える 1

1

配列のサイズを変更すると、コピー命令の数が増えるというあなたの推測は正しいです。以下に示すように、これを PTX コードとアセンブリで調べることができます。

ロード/ストア命令の最大サイズは 128 ビットです。たとえば、こちらを参照してください。つまり、あなたのケースではfloat4、の代わりにfloat.

または、プログラミング ガイドで説明されているように、データ構造の配置を明示的に指定することもできます。

#define DATABYTES 32
struct __align__(16) floatArray
{
    float a[DATABYTES/4];
};

PTX コードを表示するには、オブジェクト ファイルnvcc -c ...を生成して使用しますcubobjdump --dump-ptx objfile.o。あなたの例では、関連する部分は次のようになります。

ld.global.f32 %f1, [%rd7];
ld.global.f32 %f2, [%rd7+4];
ld.global.f32 %f3, [%rd7+8];
ld.global.f32 %f4, [%rd7+12];
ld.global.f32 %f5, [%rd7+16];
ld.global.f32 %f6, [%rd7+20];
ld.global.f32 %f7, [%rd7+24];
ld.global.f32 %f8, [%rd7+28];
ld.global.f32 %f9, [%rd7+32];
ld.global.f32 %f10, [%rd7+36];
st.global.f32 [%rd6+36], %f10;
st.global.f32 [%rd6+32], %f9;
st.global.f32 [%rd6+28], %f8;
st.global.f32 [%rd6+24], %f7;
st.global.f32 [%rd6+20], %f6;
st.global.f32 [%rd6+16], %f5;
st.global.f32 [%rd6+12], %f4;
st.global.f32 [%rd6+8], %f3;
st.global.f32 [%rd6+4], %f2;
st.global.f32 [%rd6], %f1;

配列をさらに増やすと、コンパイラがロード/ストアごとに命令を発行する代わりにループすることを選択するポイントが見つかります。

したがって、アセンブリを検査することができます cubobjdump --dump-sass objfile.o

于 2016-06-16T14:48:55.060 に答える