私は、GPU OpenCLカーネルを改良して処理を高速化することを考えています。問題は、統合されていない大量のグローバルメモリがあり、フェッチが実際にパフォーマンスを低下させていることです。だから私はグローバルメモリの多くをローカルにコピーするつもりですが、コピーするものを選ぶ必要があります。OpenCLグローバルメモリフェッチ
私の質問は、メモリの小さなチャンクの多くのフェッチは、より大きなチャンクのフェッチより多くを傷つけるのでしょうか？
clGetDeviceInfoを使用して、デバイスのキャッシュラインサイズを調べることができます。 （clGetDeviceInfo、CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE）今日の多くのデバイスでは、この値は通常16バイトです。
小さな読み込みは面倒かもしれませんが、同じキャッシュラインから読み込んでいる場合は、うまくいくはずです。短い答え：あなたは速く保つためにあなたの小さな塊を一緒に記憶に残しておく必要があります。
メモリにアクセスする2つの方法、つまりvectorAddFooおよびvectorAddBarを示すために、以下の2つの機能があります。 3番目の関数copySomeMemory（...）は、具体的にあなたの質問に適用されます。どちらのベクトル関数も作業項目を持ち、追加されるベクトルの一部を追加しますが、異なるメモリアクセスパターンを使用します。 vectorAddFooは、各作業項目を取得してベクトル要素のブロックを処理し、配列内の計算位置から開始し、ワークロードを順方向に移動します。 vectorAddBarは、次の要素をフェッチして追加する前に、gidで作業項目を開始し、gSize（=グローバルサイズ）要素をスキップします。
vectorAddBarは、読み取りと書き込みがメモリ内の同じキャッシュラインに入るため、より高速に実行されます。 4回のフロート読み取りはすべて同じキャッシュライン上にあり、メモリコントローラから実行するアクションは1つだけです。この問題で[]とb []を読み終えたら、4つの作業項目すべてが追加を行い、c []にその書き込みをキューイングできます。
vectorAddFooは、読み取りと書き込みが同じキャッシュラインにないことを保証します（非常に短いベクトル〜totalElements < 5を除く）。作業項目から読み取られるたびに、メモリコントローラからのアクションが必要になります。 gpuがすべての場合に以下の3つのフロートをキャッシュしない限り、メモリアクセスは4倍になります。
__kernel void
vectorAddFoo(__global const float * a,
__global const float * b,
__global float * c,
__global const totalElements)
{
int gid = get_global_id(0);
int elementsPerWorkItem = totalElements/get_global_size(0);
int start = elementsPerWorkItem * gid;
for(int i=0;i<elementsPerWorkItem;i++){
c[start+i] = a[start+i] + b[start+i];
}
}
__kernel void
vectorAddBar(__global const float * a,
__global const float * b,
__global float * c,
__global const totalElements)
{
int gid = get_global_id(0);
int gSize = get_global_size(0);
for(int i=gid;i<totalElements;i+=gSize){
c[i] = a[i] + b[i];
}
}
__kernel void
copySomeMemory(__global const int * src,
__global const count,
__global const position)
{
//copy 16kb of integers to local memory, starting at 'position'
int start = position + get_local_id(0);
int lSize = get_local_size(0);
__local dst[4096];
for(int i=0;i<4096;i+=lSize){
dst[start+i] = src[start+i];
}
barrier(CLK_GLOBAL_MEM_FENCE);
//use dst here...
}
一般に、より大きいサイズのフェクトがより効率的になります。私はあなたのコードを見ることなくあなたに具体的な助言を与えることはできませんが、「ストリーミング」を可能にするために作業項目から連続したチャンクにアクセスしてください。データをローカルメモリに持ってきた後に、転置やランダムなメモリアクセスを行います。
あなたは正しく質問できませんが、大域的なアクセスが大きく、再利用されている場合は、ローカルメモリを使用してください。
注：ローカルワークサイズが小さいため、データが共有されないため、使用されません。 大きなローカル作業サイズは、より少ない並列スレッドです。だからあなたは最高のものを選択する必要があります。