1

Kepler でのグローバル アトミックの実装について知りたいです。

このコードを参照してください。

1. if (threadIdx.x < workers) {
2.    temp = atomicAdd(dst, temp + rangeOffset);
3.    if (isLastPartialCalc(temp)) {                            
4.        atomicAdd(dst,-300000.0f);
5.    }
6. }

このために4行目を変更すると：

*dst -= 300000.0f;

性能が落ちる！これ以上スレッドがこの値に書き込むことはないため、変更は安全です (出力は同じです)。

アトミックを使用するカーネル: ~883us gmem を直接使用するカーネル: ~903us

私は何度か実行しましたが、常に変更に対して約 20us のペナルティが発生します。

更新 アトミックを使用しないストアは常にL2でミスを生成するようですが、アトミックバージョンは常にヒットを生成します...したがって、「アトミック」でフラグが立てられた(または何か)場所に書き込もうとすると、 L2 であり、gmem に別のリクエストを行います

4

2 に答える 2

1

このキャッシュ ラインの更新は、(特定のコードでは) 2 回目のグローバル アトミック アクセスより明らかにコストがかかります。

1 つの SM から Kepler GK110 (K20 など) 上のグローバル メモリへの 1 つのグローバル アトミック アクセスは、実際には非常に高速です。

ケプラー ホワイト ペーパーに示されているように、ケプラーはフェルミに比べてグローバル アトミックの速度が向上しています。

共通のグローバル メモリ アドレスに対するアトミック操作のスループットは、1 クロックあたり 1 つの操作に 9 倍向上します。

于 2013-07-01T13:12:15.233 に答える