は、次のコードを考えてみましょう:共有メモリ配列のデフォルト値を設定する方法はありますか?
__global__ void kernel(int *something) {
extern __shared__ int shared_array[];
// Some operations on shared_array here.
}
は、それはいくつかの値に全体のshared_arrayを設定することが可能です - 例えば0 - 各スレッドを明示的に指定せずに
は、次のコードを考えてみましょう:共有メモリ配列のデフォルト値を設定する方法はありますか?
__global__ void kernel(int *something) {
extern __shared__ int shared_array[];
// Some operations on shared_array here.
}
は、それはいくつかの値に全体のshared_arrayを設定することが可能です - 例えば0 - 各スレッドを明示的に指定せずに
いいえ、共有メモリは初期化されていません。あなたは、CUDA Cプログラミング・ガイド3.2、セクションB.2.4.2から
...何とか、それを一つの方法または別の自分自身を初期化する必要があり、段落2:
__shared__
変数は一部として初期化を持つことができませんその宣言の
これは、共有変数の重要なデフォルトコンストラクタも破棄します。
あなたが効率的にこの
// if SHARED_SIZE == blockDim.x, eliminate this loop
for (int i = threadIdx.x; i < SHARED_SIZE; i += blockDim.x)
shared_array[i] = INITIAL_VALUE;
__syncthreads();
のような並列に共有配列を初期化することができますはい、できます。あなたは、他のは、例えばないながらブロックの最初のスレッドは、それを設定することを指定することができます。:
extern __shared__ unsigned int local_bin[]; // Size specified in kernel call
if (threadIdx.x == 0) // Wipe on first thread - include " && threadIdx.y == 0" and " && threadIdx.z == 0" if threadblock has 2 or 3 dimensions instead of 1.
{
// For-loop to set all local_bin array indexes to specified value here - note you cannot use cudaMemset as it translates to a kernel call itself
}
// Do stuff unrelated to local_bin here
__syncthreads(); // To make sure the memset above has completed before other threads start writing values to local_bin.
// Do stuff to local_bin here
syncthreadsを呼び出す前に、これは他のすべてを可能として、理想的には、できるだけ多くの作業を行う必要がありますスレッドがmemsetが完了する前に作業を行うようにします。例えば、条件付き分岐がある場合など、作業がかなり異なるスレッド完了時間を持つ可能性がある場合にのみ、これは重要です。 スレッド0のfor-loopを設定する場合、カーネルのパラメータとしてlocal_bin配列のサイズを渡す必要があるので、反復処理する配列のサイズを知る必要があります。あなたはもちろん、1Dのブロックを持っている場合にのみケースだ
ありがとうございます。私は最終的な実装でこれに類似したものを使用しました。 – rayryeng
これは並列化の利点を失い、常にthreadIdxとBlockIdxを可能な限り使用しようとします – ejectamenta
私はあなたが一般的だと思いますが、それらから初期化値を計算できるかどうか、/etcの割り当てパターン。 – metamorphosis
。どんな初心者も明白なトラップにはならないと言っています。私はまた、浮動小数点数をどれくらい後押しするのか疑問に思います。これは、新しいデバイスではまだまだ別のトリックですが、このメモリ統合型のinitとどのくらいのメリットがありますか。 Sidenote、2Dまたは3Dカーネルの中に読み込んでいる場合、[z] [y] [x]の配列のようなワープに分割されていることを知ることはイントラントです。だから、[x]スレッドはお互いに最も近く、そして[x]スレッドは最も遠くに書くようにしてください。 –
私は試してみましたが、reinterpret_castを使ってhttps://devblogs.nvidia.com/parallelforall/cuda-pro-tip-increase-performance-with-vectorized-memory-accessのような16バイトのチャンクでコピーしています/まだ少し良い結果が得られます。また、コピーしたいデータに__restrict__を設定することも重要です。簡単に+ 15%のパフォーマンスを向上させることができます。これは、16バイトの整列チャンクほどです。 –
メモリの統合は、16倍の速度差を容易に意味することに注意してください。 –