これを行う1つの方法は、グローバルフラグ変数を使用することです。値を見つけたら、アトミックに1に設定します。他のスレッドは、作業中にその値をチェックします。
例えば:
__kernel test(__global int* buffer, __global volatile int* flag)
{
int tid = get_global_id(0);
int sx = get_global_size(0);
int i = tid;
while(buffer[i] != 8) //Whatever value we're trying to find.
{
int stop = atomic_add(&flag, 0); //Read the atomic value
if(stop)
break;
i = i + sx;
}
atomic_xchg(&flag, 1); //Set the atomic value
}
これにより、カーネル全体を実行するよりもオーバーヘッドが増える可能性があります(すべての反復で多くの作業を行っている場合を除く)。さらに、各スレッドが配列内の単一の値をチェックしているだけの場合、このメソッドは機能しません。各スレッドには、複数回の作業の反復が必要です。
最後に、アトミック変数への書き込みがすぐにコミットされない場合を確認しました。そのため、書き込みがコミットされていないために、このコードがシステムでデッドロックするかどうかを確認する必要があります。