0

値の大きな配列 (まだ 64 kB より小さい) があるとします。これはカーネルで非常に頻繁に読み取られますが、書き込まれることはありません。ただし、外部から変更することはできます。配列には 2 つの値のセットがあり、それらを left と right と呼びましょう。問題は、大きな配列を __global として取得し、それを __local left および __local right 配列に書き込む方が高速かどうかです。または定数 __constant として取得し、カーネルでアクセスを処理しますか? 例えば：

__kernel void f(__global large, __local left, __local right, __global x, __global y) {
    for(int i; i < size; i++) {
        left[i] = large[i];
        right[i] = large[i + offset];
    }
    ...
    x = foo * left[idx];
    y = bar * right[idx];
}

対:

__kernel void f(__constant large, __global x, __global y) {
    ...
    x = foo * large[idx];
    y = bar * large[idx * offset];
}

(インデックス作成はもう少し複雑ですが、たとえば、マクロを使用して作成できます) 定数メモリはグローバル空間にあると読みましたが、遅くする必要がありますか? Nvidia カードで動作します。

4

1 に答える 1

1

まず、2番目のケースでは、CPUで結果を利用できるようにする方法が必要です。global space計算後にコピーして戻すことを想定しています。

カーネルで何をするかによると思います。たとえば、カーネルの計算が重い場合（スレッドごとに多くの計算が行われる場合）、最初のオプションで十分になる可能性があります。なんで？

  • スペースからglobal largeスペースにデータをコピーするのにある程度の時間を費やします。localleftright
  • ローカル空間のデータに対して多くの計算を行います-OK
  • local leftからとrightへのコピーバックに時間を費やしますglobal large。- 許容できる。

ただし、カーネルが比較的軽い場合、つまり各スレッドがいくつかの小さな計算を行う場合は、

  • 宇宙のデータを使っていくつかの計算を行いますconstant。これはおそらく、あまりアクセスする必要がないことを意味します。
  • 中間結果をローカルスペースに保存します。
  • local宇宙から宇宙へコピーバックするのに少し時間を費やしglobalます。- 許容できる。

大規模なカーネルの場合、最初のオプションの方が適しています。小さなカーネルの場合は2番目。

PSもう1つ注意すべき点は、次々に動作する複数のカーネルがある場合はlarge、必ず最初のオプションを選択することです。これにより、データをグローバルメモリスペースに保持でき、カーネルを起動するたびにコピーを実行する必要がなくなります。

編集：あなたはそれが非常に頻繁にアクセスされると言ったので、私はあなたがおそらく最初のオプションで行くべきだと思います。

于 2013-03-06T21:50:55.583 に答える