2

CUDA Cプログラミングガイド（p.70）には、次のように書かれています。

グローバルメモリはデバイスメモリに常駐し、デバイスメモリは32、64、または128バイトのメモリトランザクションを介してアクセスされます。これらのメモリトランザクションは自然に整列する必要があります。サイズに整列された（つまり、最初のアドレスがサイズの倍数である）デバイスメモリの32、64、または128バイトセグメントのみがメモリトランザクションによって読み取りまたは書き込みできます。 。

それで、デバイス機能で一度に32、64、または128バイトにアクセスしたい場合（たとえば、共有メモリへのコピーの場合）、この操作に最も適した機能（または割り当て）は何ですか？

従来のCmemcpy関数は、一度に32バイトにアクセスしないようです（非常に遅いです）。また、これはベクターデータではないため、1つのスレッドで一度にこのデータを読み取ってほしいと思います。

dbauppに

memcpyはうまく機能しますが、私は速度について話しています。たとえば、デバイスメモリポインタpがあり、デバイス関数で次のコードを実行するとします。

a）char c [8]; memcpy（c、p、8）;

b）char c [8]; *（double *）c = *（double *）p;

上記の2つのケースの場合、結果は同じですが、ケースbはケースaよりもほぼ8倍高速です（コードでテストおよび確認しました）。

また、参考までに、cudaMemcpy関数はデバイス関数では機能しません。

だから、私が知りたいのは、1回の操作で16バイトをコピーする方法があるかどうかです。（うまくいけば、memcpy（c、p、16）;より16倍高速です）;

4

1 に答える 1

6

あなたが何をしようとしているのかは100％明確ではありません。グローバルメモリから共有メモリにデータをコピーしようとしている場合、おそらくそれは何らかの構造を持っています。たとえば、charsやfloatsの配列などです。次の回答は、sの配列を処理していることを前提としています（任意のデータ型charに置き換えることができます）。char

要約：一度に32/64/128バイトに明示的にアクセスすることは考えずに、メモリアクセスを合体できるようにコードを記述してください。

通常のC/C ++などと同じように、CUDAを使用してデータにアクセスできます。個々のバイトにまで到達することもできます。プログラミングガイドが言っていることは、データにアクセスするときはいつでも、32/64/128バイトのブロックを読み取る必要があるということです。たとえば、char a[128]取得したい場合a[17]、GPUはからデータを取得できるようにするためにa[0]から読み取る必要があります。これは透過的に行われます。たとえば、個々のバイトにアクセスできるようにするために、別の方法でコーディングする必要はありません。a[31]a[17]

主な考慮事項はメモリアクセス速度です。情報バイトごとに31ジャンクバイトを読み取る必要がある場合、実効メモリ帯域幅を32分の1に削減します（つまり、さらに多くのグローバルメモリアクセスを実行する必要があります。 sloowww）！

ただし、GPUのメモリアクセスは、ブロック内のスレッド間で「合体」することができます（この質問は、合体を最適化するための妥当な開始点を提供します）。簡単に言うと、合体により、ブロック内の複数のスレッドで同時に発生するメモリアクセスを「バッチ処理」できるため、1回の読み取りのみを実行する必要があります。

これのポイントは、（単一のスレッド内ではなく）ブロック内のスレッド間で合体が発生することです。したがって、共有メモリへのコピーの場合は、次のことができます（グローバルメモリ内arrayのsの配列）。char

__shared__ char shrd[SIZE];

shrd[threadIdx.x] = array[blockDim.x * blockIdx.x  + threadIdx.x];
__syncthreads();

これにより、各スレッドが1バイトを共有配列にコピーします。このmemcpy操作は基本的に並行して行われ、データアクセスは合体されるため、無駄な帯域幅（または時間）はありません。

上記の戦略は、単一のスレッドを反復処理してバイトごとにコピーするよりもはるかに優れています。

配列のnバイトの各ブロックを単一のnバイトのデータ型として扱い、各スレッドにそれをコピーさせることもできます。たとえば、n == 16の場合、いくつかのキャストを実行しますuint4

__shared__ char shrd[SIZE];

((uint4*)shrd)[threadIdx.x] = ((uint4*)array)[blockDim.x * blockIdx.x  + threadIdx.x];
__syncthreads();

これにより、各スレッドが一度に16バイトをコピーできるようになります。コードのそのビットに関する注記：

  • 私はそれをテストまたはベンチマークしていません
  • それが良い習慣かどうかはわかりません（そうでないことを強く期待します）。
  • インデックスは16でスケーリングされます（たとえばthreadIdx.x == 1、への書き込みに対応しますshrd[16],shrd[17],...,shrd[31]）

補足として：特定のユースケースによっては、組み込みのcudaMemcpy関数が役立つ場合があります。

于 2012-05-31T14:29:02.913 に答える