次のように動作するワークフローがあります。
- 初期値のロード
- 中間結果へのプロセス値A
- プロセスAから中間結果B
- プロセスBから中間結果C
- プロセスCおよびBから中間結果DおよびE
- 部分Dを合計して最終結果F
すべての中間結果の自然な構造は、cudaMallocPitch()を使用して割り当てている2D配列の構造です。
残念ながら、私のアルゴリズムでは、D、E、C、およびBを一度にメモリに保持する必要があり、D&Eは、個別に、Bよりもメモリ内で4倍大きくなります。処理に別の制限があるため(グラフ構造を反復処理)メモリ内)、AまたはBの次元は、DとEの最大次元によって制限されます。これらは、初期値のメモリ使用量+Bのメモリ消費量+Cのメモリ消費量によって決定されます。この依存関係は、(非常に大きな問題セットに対応するために)ホストからデバイスメモリとの間で中間結果のセクションを「ページング」しているためです。ステップ1〜3が全体で完了するまで、ステップ4を開始できません。問題セット。
問題セット全体に対してBを取得したら、Aを削除できます。
私は現在、次の機能を使用してD+Eの最大サイズを決定しています。
int gpuCalculateSimulPatterns(int lines, int patterns) {
// get free memory
size_t free_mem, total_mem;
int allowed_patterns;
cudaMemGetInfo(&free_mem, &total_mem);
allowed_patterns = (free_mem - (lines*sizeof(int))) / (lines*(sizeof(int)*2.5) + lines*sizeof(char)*1.5);
return min(patterns, allowed_patterns -(allowed_patterns % 32));
}
これは「機能」しますが、DまたはEのサイズ(サイズとメモリ使用量は同じ)を25%過大評価し、予想されるBのサイズを2倍にしているためです。それでも、メモリが不足しているエッジケースに遭遇します。メモリが不足しているため、割り当ては失敗します。私のカーネルはグローバルメモリとの間で複数の読み取りと書き込みを行うため、カード上のメモリをより効率的に使用し、アライメントを維持したいと考えています。
いいえ、共有メモリを使用することはできません。複数のブロックにまたがって複数のカーネルを使用しており、ブロック内のスレッドはまったく相互作用しないためです。
cudaMallocPitch()は、正常に割り当てられたメモリの使用済みピッチのみを返すことがわかりました。ドライバーに2Dメモリ割り当て要求を渡して、割り当てられるピッチを要求する方法はありますか?
試行/エラー最適化ルーチンを作成しますが、A、B、D、およびE間のディメンションのリンクされた依存関係(CIはピッチ線形に割り当てられていないため、事前に計算します)により、これはお粗末なソリューションになり、必要になります問題セットごとに再計算されます。
任意の量のデバイスメモリに収まる中間データセットの適切なサイズを決定できる、より良いアプローチはありますか?
編集:
中間Aのメモリが再利用されています。私の境界計算では、C + D + E +B>>初期+A+ B(AとBが1バイトの文字であるという事実により当てはまります)と仮定しています。同じ寸法ですが、C、D、Eはintです)など、B + C + D+Eに十分なスペースがあることを確認するだけで済みます。
これをテストするためにComputeCapability2.xカードのみを使用しています(Quadro 2000、Tesla C2075、GTX460)。