CUDA で連立方程式を解くアプリケーションがあります。各スレッドが最大 4 つの解を見つけることができることは確かですが、ホストにコピーして戻すにはどうすればよいですか?
すべてのスレッドが 4 つのソリューション (ソリューションごとに 4 つの double) を格納するのに十分なスペースを持つ巨大な配列と、スレッドごとのソリューションの数を持つ別の配列を渡していますが、これは単純なソリューションであり、カーネルの現在のボトルネックです。
私はこれを最適化するのが本当に好きです。主な問題は、スレッドごとに可変数のソリューションを単一の配列に連結することです。
あなたが探している機能は、ストリーム圧縮と呼ばれます。
結果をコンパクトな形式で直接保存しようとすると、スレッド間に非常に多くの依存関係が作成される可能性が高いため、スレッドごとに 4 つのソリューションの余地を含む配列を提供する必要がある可能性があり、より少ないデータをコピーして戻すことができるため、パフォーマンスが向上します。カーネルの実行時間が長くなると、ホストが失われます。これに対する例外は、ほとんどすべてのスレッドで解決策が見つからない場合です。その場合、アトミック操作を使用して配列へのインデックスを維持できる場合があります。したがって、見つかったソリューションごとに、配列のインデックスに格納し、アトミック操作を使用してインデックスを増やします。これにはatomicAdd()を使用するのが安全だと思います。結果を格納する前に、スレッドは atomicAdd() を使用してインデックスを 1 つ増やします。atomicAdd() は古い値を返します。
ただし、かなりの数の結果が得られる、より一般的な状況を考えると、最適な解決策は、別の手順として圧縮操作を実行することです。これを行う 1 つの方法は、 を使用することですthrust::copy_if。背景については、この質問を参照してください。