ポインタでasync_work_group_copy（）を使用していますか？

Question

__kernel void kmp(__global char pattern[1*4], __global char* string, __global int failure[1*4], __global int ret[1], int g_length, int l_length, int thread_num){ 
    int pattern_num = 1; 
    int pattern_size = 4; 
    int gid = get_group_id(0); 
    int glid = get_global_id(0); 
    int lid = get_local_id(0); 
    int i, j, x = 0; 

    __local char *tmp_string; 
    event_t event; 

    if(l_length < pattern_size){ 
      return; 
    } 

    event = async_work_group_copy(tmp_string, string+gid*g_length, g_length, 0); 
    wait_group_events(1, &event); 

これは私のコードの一部です。

私はテキストに一致するパターンを見つけたいと思います。

最初に、CPU側ですべてのパターンと文字列（テキストから文字列を読み取り、実験的に1つのパターンのみを使用）を初期化します。

次に、それらをkmpという名前のカーネルに転送します。 （パラメータl_lengthとg_lengthはlidにコピーされ、それぞれglidになる文字列のサイズです。つまり、文字列）

最後に、分割された文字列をローカルメモリにコピーします。

しかし、問題があります。 async_work_group_copy（）を使用してコピーすると、有効な結果が得られません。

__local char * tmp_stringを配列に変更すると、問題は残ります。

私がしたいのは、1）文字列を分割し、2）各スレッドにコピーして3）、一致する数値を計算します。

このコードで何が問題なのでしょうか。ありがとう！

Answer 1

のOpenCLの仕様では、これを持っている：

非同期コピーがワークグループ内のすべての作業項目によって実行され、この 組み込み関数は、したがって、 仕事にすべての作業項目が遭遇しなければなりません-group同じ引数値でカーネルを実行する。 それ以外の場合は結果が未定義です。

したがって、グループ内の作業項目の早期復帰は避けてください。とにかく早期復帰がCPUに適しています。これがGPUの場合は、拡張/パディングされた入出力バッファを使用して最後のオーバーフロー部分を計算するだけです。

それ以外の場合は、デバイスが作業項目（専用の秘密のパイプライン）を使用していない限り、グループ全体を早期に返すことができます（作業項目が非同期のコピー命令に当たらないため動作します）非同期コピー操作の場合

たぶん、あなたはworkgroupsize = remaining_sizeの代わりに、余分なバッファサイズや制御ロジックを持つ最新のアイテムを残り計算するために（同時に別のキューに）は、第2のカーネルをエンキューすることができます。

---

tmp_stringあなたはそれから/に何かをコピーしようとしている場合、割り当てられた/初期化する必要があります。だから、おそらくあなたはそれの配列バージョンが必要になります。

---

async_work_group_copyそれがグローバルに非同期コピーに使用するローカルメモリの最新のビットを取得する前に、その障壁を必要とする同期ポイントではありません。

  __kernel void foo(__global int *a, __global int *b) 
      { 
       int i=get_global_id(0); 
       int g=get_group_id(0); 
       int l=get_local_id(0); 
       int gs=get_local_size(0); 
       __local int tmp[256]; 
       event_t evt=async_work_group_copy(tmp,&a[g*gs],gs,0); 
       // compute foobar here in async to copies 
       wait_group_events(1,&evt); 

       tmp[l]=tmp[l]+3; // compute foobar2 using local memory 
       barrier(CLK_LOCAL_MEM_FENCE); 

       event_t evt2=async_work_group_copy(&b[g*gs],tmp,gs,0); 
       // compute foobar3 here in async to copies 
       wait_group_events(1,&evt2); 
      }