4

2 つの事実: CUDA 5.0 では、後でリンクするために CUDA コードをさまざまなオブジェクト ファイルにコンパイルできます。CUDA アーキテクチャ 2.x では、関数が自動的にインライン化されなくなりました。

C/C++ ではいつものように、関数__device__ int foo()を に実装し、functions.cuそのヘッダーを に配置しましたfunctions.hu。関数fooは、他の CUDA ソース ファイルで呼び出されます。

を調べると、それがローカル メモリに流出してfunctions.ptxいることがわかります。foo()テストの目的で、私は の肉のすべてをコメントし、foo()それreturn 1; を作成しました.ptx. （関数は何もしないので、それが何であるか想像できません！）

しかし、実装をfoo()ヘッダー ファイル に移動して修飾子functions.hu を追加すると、__forceinline__ローカル メモリには何も書き込まれません。

ここで何が起こっているのですか？ CUDA がこのような単純な関数を自動的にインライン化しないのはなぜですか?

個別のヘッダーと実装ファイルの全体的なポイントは、コードの保守を容易にすることです。しかし、ヘッダーとそれらに一連の関数 (またはそれらすべて) を貼り付ける必要がある場合は__forceinline__、CUDA 5.0 のさまざまなコンパイル ユニットの目的を無効にします...

これを回避する方法はありますか？

シンプルで実際の例:

functions.cu:

__device__  int  foo
        (const uchar param0,
        const uchar *const param1,
        const unsigned short int param2,
        const unsigned short int param3,
        const uchar param4) 
{    
    return 1; //real code commented out.
}

上記の関数は、ローカル メモリにスピルします。

functions.ptx:

.visible .func  (.param .b32 func_retval0) _Z45fooPKhth(
        .param .b32 _Z45foohPKhth_param_0,
        .param .b64 _Z45foohPKhth_param_1,
        .param .b32 _Z45foohPKhth_param_2,
        .param .b32 _Z45foohPKhth_param_3
)
{
        .local .align 8 .b8     __local_depot72[24];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .s16       %rc<3>;
        .reg .s16       %rs<4>;
        .reg .s32       %r<2>;
        .reg .s64       %rd<2>;
4

1 に答える 1

4

すべてのローカル メモリ使用量がスピルを表しているわけではありません。呼び出される関数は、ローカル メモリにあるスタック フレームの作成を含む ABI 呼び出し規則に従う必要があります。nvcc にコマンドライン スイッチ -Xptxas -v が渡されると、コンパイラはスタックの使用状況とスピルをそのサブコンポーネントとして報告します。

現在 (CUDA 5.0)、CUDA ツールチェーンは、一部のホスト コンパイラのように、コンパイル ユニットの境界を越えた関数のインライン化をサポートしていません。したがって、個別のコンパイルの柔軟性 (長いコンパイル時間で大規模なプロジェクトの小さな部分のみを再コンパイルすることや、デバイス側のライブラリを作成する可能性など) と、通常は関数から生じるパフォーマンスの向上との間にはトレードオフがあります。インライン化 (たとえば、ABI 呼び出し規約によるオーバーヘッドの排除、関数境界を越えた定数伝搬などの追加の最適化を可能にする)。

単一のコンパイル ユニット内での関数のインライン展開は、コンパイラのヒューリスティックによって制御されます。コンパイラのヒューリスティックは、インライン展開がパフォーマンスの点で有益であるかどうか (可能であれば) を判断しようとします。これは、すべての関数がインライン化されているわけではないことを意味します。プログラマーは、関数属性__forcinline__およびを使用してヒューリスティックをオーバーライドできます__noinline__。

于 2013-06-14T16:49:12.743 に答える