8

ではOpenCL、バッファにREAD_ONLYまたはのフラグを付けるとパフォーマンス上の利点はありますかWRITE_ONLY?

これkernelは私がよく目にするものです (a はREAD_ONLYb ですWRITE_ONLY):

__kernel void two_buffer_double(__global float* a, __global float* b)
{
    int i = get_global_id(0);
    b[i] = a[i] * 2;
}

kernelグローバルメモリの使用量が少ないため、これはより良いようです（ a is READ_WRITE）：

__kernel void one_buffer_double(__global float* a)
{
    int i = get_global_id(0);
    a[i] = a[i] * 2;
}

READ_ONLYおよびWRITE_ONLYフラグは、デバッグとエラーの検出を支援するために存在しますか?

4

3 に答える 3

7

実際には2種類あることに注意してください。バッファを割り当てるときにとがありますがCL_MEM_READ_ONLY、カーネル コードでポインタを装飾するためにともあります。CL_MEM_WRITE_ONLYCL_MEM_READ_WRITE__read_only__write_only__read_write

これらは、最適化とエラー チェックの両方に使用できます。まずは性能を見てみましょう。書き込み専用バッファが検出された場合、書き込みをキャッシュする必要がなく (ライト スルー キャッシュのように)、読み取り用により多くのキャッシュを節約できます。これは GPU ハードウェアに大きく依存し、少なくとも NVIDIA ハードウェアには、これを実際に実装するために必要な命令 (.csおよび.lu修飾子) があります。PTX ISAを参照できます。コンパイラが実際にこの最適化を実行しているという証拠は見たことがありません。

__kernel void Memset4(__global __write_only unsigned int *p_dest,
    const unsigned int n_dword_num)
{
    unsigned int i = get_global_id(0);
    if(i < n_dword_num)
        p_dest[i] = 0; // this
}

次のようにコンパイルされます。

st.global.u32 [%r10], %r11; // no cache operation specified

CUDA にはこれらの修飾子に相当するものがないため、これは理にかなっているため、コンパイラはそれらを黙って無視する可能性が最も高いです。しかし、それらをそこに置いても害はありません。将来的にはもっと幸運になるかもしれません. CUDA では、__ldg関数を使用し、コンパイラ フラグを使用して、L1 でのグローバル メモリ転送のキャッシュをオプトイン/オプトアウトすることで、この機能の一部が公開されます ( -Xptxas -dlcm=cg)。asmキャッシュをバイパスすると大きな利点が得られる場合は、いつでも使用できます。

constエラーチェックに関しては、カーネル宣言で指定子を使用して、読み取り専用バッファーへの書き込みを簡単に回避できます。純粋な「C」では、書き込み専用バッファからの読み取りを禁止することはできません。

これらのバッファをホストメモリにマッピングするときに、別の最適化が行われる可能性があります。バッファをマッピングするときCL_MEM_READ_ONLY、ホストはそのメモリに書き込むだけで、デバイスはそれを読み取るだけなので、マッピングされた領域は初期化されていないままになる可能性があります。同様に、CL_MEM_WRITE_ONLYバッファーのマッピングを解除する場合、ドライバーは (ホストによって変更される可能性がある) 内容をホスト メモリからデバイス メモリにコピーする必要はありません。私はこれを測定しませんでした。

補足として、私は使用してみました：

inline unsigned int n_StreamingLoad(__global __read_only const unsigned int *p_src)
{
#ifdef NVIDIA
    unsigned int n_result;
    asm("ld.global.cs.u32 %r0, [%r1];" : "=r" (n_result) : "r" (p_src));
    return n_result;
#else // NVIDIA
    return *p_src; // generic
#endif // NVIDIA
}

inline void StreamingWrite(__global __write_only unsigned int *p_dest, const unsigned int n_value)
{
#ifdef NVIDIA
    asm("st.global.cs.u32 [%r0], %r1;" : : "r" (p_dest), "r" (n_value) : "memory");
#else // NVIDIA
    *p_dest = n_value; // generic
#endif // NVIDIA
}

sm_35これにより、デバイスを備えた単純な memcpy カーネル (GTX 780 および K40 でテスト済み)でも、約 15 GB/秒余分に使用できます。目立ったスピードアップは見られませんでしたsm_30(そこでサポートされることを意味するかどうかはわかりませんが、命令は ptx から削除されていません)。自分で定義する必要があることに注意してNVIDIAください (または、カーネル コードで OpenCL デバイス ベンダーを検出する を参照してください)。

于 2015-12-10T14:30:49.367 に答える
5

場合によります、

READ_ONLY __globalメモリの場所は「グローバル/コンスタント メモリ データ キャッシュ」に格納されます。これは通常のキャッシュや GPU の RAM (こちらを参照) よりもはるかに高速ですが、CPU では問題ありません。

WRITE_ONLY の利点はわかりません。GPU は、キャッシュを必要とせずにデータをストリーミングできることを認識しているため、それも役立つかもしれません。

わからない場合は、行って測定してください...

于 2013-07-27T19:31:07.843 に答える
5

あなたの質問に率直に答えるには、次のように言います。いいえ、これらのフラグは、デバッグやエラーの検出を支援するためだけに存在するわけではありません。ただし、これらのフラグが実装でどのように使用され、パフォーマンスにどのように影響するかについての参照を提供することは困難です。

私の理解では(残念ながら、どのドキュメントにも裏付けられていません)、これらのフラグを使用すると、バッファーの使用方法にさらに制約を課すことができるため、ランタイム/ドライバー/コンパイラーがパフォーマンスを向上させる可能性のあるいくつかの仮定を立てるのに役立ちます。たとえば、作業項目は書き込みを想定していないため、カーネルがそれを使用している間、読み取り専用バッファーとのメモリの一貫性について心配する必要はないと思います。したがって、いくつかのチェックがスキップされる可能性があります... ただし、Opencl では、バリアなどを使用して自分でこれを処理する必要があります。

また、Opencl 1.2 以降、ホストがバッファにアクセスする必要がある方法に関連するいくつかの他のフラグが今回導入されたことにも注意してください。がある：

CL_MEM_HOST_NO_ACCESS,
CL_MEM_HOST_{READ, WRITE}_ONLY,
CL_MEM_{USE, ALLOC, COPY}_HOST_PTR

これもまた、opencl を実装する人々がパフォーマンスを向上させるのに役立つはずだと推測していますが、AMD または NVIDIA の専門家からの意見が必要になると思います。

これまでに述べたことはすべて私の考えであり、深刻なドキュメントに基づいていないことに注意してください(私は何も見つけることができませんでした)。

一方、@Quonuxが述べたように、標準が読み取り専用バッファーを一定のスペースに配置することを強制していないことは確かです。一部の実装では、小さなバッファーに対してこれを行う場合があります。定数空間メモリは小さいため、読み取り専用バッファが大きすぎて収まらないことを忘れないでください。バッファが定数空間メモリにあることを確認する唯一の方法は、カーネル コードで定数キーワードを次のように使用することです。ここで説明します。もちろん、ホスト側で定数バッファーを使用する場合は、読み取り専用フラグを使用する必要があります。

于 2013-07-27T21:21:36.897 に答える