と を使用して線形メモリを割り当ててコピーする方法については、かなりよく理解していcudaMalloc()ますcudaMemcpy()。しかし、CUDA 関数を使用して 2D または 3D 行列を割り当ててコピーする場合、特に 2D/3D 配列を扱うときに常に存在するピッチ付きポインターに関して、さまざまな引数に困惑することがよくあります。ドキュメントは、それらの使用方法に関するいくつかの例を提供するのに適していますが、パディングとピッチの概念に精通していることを前提としていますが、私はそうではありません.
私は通常、ドキュメントやウェブ上のどこかで見つけたさまざまな例を微調整することになりますが、それに続くブラインドデバッグは非常に苦痛なので、私の質問は次のとおりです。
ピッチとは？どうやって使うの？CUDA で 2D および 3D 配列を割り当ててコピーするにはどうすればよいですか?
ここでは、CUDA のピッチド ポインターとパディングについて説明します。
まず、非線形メモリが存在する理由から始めましょう。cudaMalloc でメモリを割り当てると、結果は malloc での割り当てのようになります。指定されたサイズの連続したメモリ チャンクがあり、そこに必要なものを何でも入れることができます。10000 float のベクトルを割り当てたい場合は、単純に次のようにします。
float* myVector;
cudaMalloc(&myVector, 10000*sizeof(float));
次に、従来のインデックス作成によって myVector の i 番目の要素にアクセスします。
float element = myVector[i];
次の要素にアクセスしたい場合は、次のようにします。
float next_element = myvector[i+1];
最初の要素のすぐ隣にある要素にアクセスするのは（私が気付いていない理由で、今のところしたくない理由で）安価であるため、非常にうまく機能します。
メモリを 2D 配列として使用すると、状況が少し異なります。10000 float ベクトルが実際には 100x100 配列であるとしましょう。同じ cudaMalloc 関数を使用して割り当てることができます。i 番目の行を読み取りたい場合は、次のようにします。
float* myArray;
cudaMalloc(&myArray, 10000*sizeof(float));
int row[100]; // number of columns
for (int j=0; j<100; ++j)
row[j] = myArray[i*100+j];
したがって、メモリを myArray+100*i から myArray+101*i-1 に読み取る必要があります。必要なメモリ アクセス操作の数は、この行が必要とするメモリ ワードの数によって異なります。メモリ ワードのバイト数は、実装によって異なります。1 つの行を読み取るときのメモリ アクセスの数を最小限に抑えるには、単語の先頭で行を開始することを保証する必要があります。したがって、新しい行の開始まですべての行のメモリをパディングする必要があります。
配列をパディングするもう 1 つの理由は、共有メモリ アクセスに関する CUDA のバンク メカニズムです。配列が共有メモリにある場合、配列は複数のメモリ バンクに分割されます。同じメモリ バンクに属するメモリにアクセスしない限り、2 つの CUDA スレッドが同時にアクセスできます。通常、各行を並行して処理する必要があるため、各行を新しいバンクの先頭にパディングすることで、シミュレートしてアクセスできるようにすることができます。
ここで、cudaMalloc で 2D 配列を割り当てる代わりに、cudaMallocPitch を使用します。
size_t pitch;
float* myArray;
cudaMallocPitch(&myArray, &pitch, 100*sizeof(float), 100); // width in bytes by height
ここでのピッチは関数の戻り値であることに注意してください: cudaMallocPitch は、システム上でどうあるべきかをチェックし、適切な値を返します。cudaMallocPitch が行うことは次のとおりです。
最後に、各行が のサイズではなくピッチのサイズになっているため、通常、必要以上に多くのメモリを割り当てていますw*sizeof(float)。
しかし、列内の要素にアクセスしたい場合は、次のようにする必要があります。
float* row_start = (float*)((char*)myArray + row * pitch);
float column_element = row_start[column];
2 つの連続する列の間のバイト単位のオフセットは、配列のサイズから推測できなくなりました。そのため、cudaMallocPitch によって返されるピッチを維持したいと考えています。また、ピッチはパディング サイズ (通常はワード サイズとバンク サイズの最大値) の倍数であるため、うまく機能します。わーい。
cudaMallocPitch によって作成された配列内の単一の要素を作成してアクセスする方法がわかったので、線形であるかどうかにかかわらず、他のメモリとの間でその要素全体をコピーしたい場合があります。
malloc を使用してホストに割り当てられた 100x100 配列に配列をコピーするとします。
float* host_memory = (float*)malloc(100*100*sizeof(float));
cudaMemcpy を使用する場合、cudaMallocPitch で割り当てられたすべてのメモリを、各行間のパディングされたバイトを含めてコピーします。メモリのパディングを避けるために必要なことは、各行を 1 つずつコピーすることです。手動で行うことができます：
for (size_t i=0; i<100; ++i) {
cudaMemcpy(host_memory[i*100], myArray[pitch*i],
100*sizeof(float), cudaMemcpyDeviceToHost);
}
または、利便性のためにパディング バイトを割り当てたメモリから有用なメモリのみが必要であることを CUDA API に伝えることができます。そのため、独自の混乱を自動的に処理できれば、非常に便利です。ありがとうございます。cudaMemcpy2D に入ります。
cudaMemcpy2D(host_memory, 100*sizeof(float)/*no pitch on host*/,
myArray, pitch/*CUDA pitch*/,
100*sizeof(float)/*width in bytes*/, 100/*heigth*/,
cudaMemcpyDeviceToHost);
これで、コピーは自動的に行われます。width（ここでは100xsizeof(float)）、height time（ここでは100）、スキップピッチで指定されたバイト数をコピーします次の行にジャンプするたびにバイト。宛先メモリもパディングされる可能性があるため、宛先メモリのピッチを提供する必要があることに注意してください。ここではそうではないので、ピッチはパディングされていない配列のピッチと同じです。これは行のサイズです。memcpy 関数の幅パラメーターはバイト単位で表されますが、高さパラメーターは要素数で表されることにも注意してください。これは、上記の手動コピーを書いたように、コピーが行われる方法によるものです。幅は、行に沿った各コピーのサイズ (メモリ内で連続する要素) であり、高さは、この操作が必要な回数です。達成される。(これらの単位の不一致は、物理学者として、私を非常に悩ませます。)
3D 配列は実際には 2D 配列と変わらず、追加のパディングは含まれていません。3D 配列は、パディングされた行の単なる 2D の古典的な配列です。そのため、3D 配列を割り当てるときに、行に沿った連続するポイント間のバイト カウントの差である 1 つのピッチのみが取得されます。深さの次元に沿って連続するポイントにアクセスする場合は、ピッチに列の数を安全に掛けて、slicePitch を得ることができます。
3D メモリにアクセスするための CUDA API は、2D メモリ用のものとは少し異なりますが、考え方は同じです。