2.19. メモリストライド

メモリストライドはキャッシュ等の限られた幅の領域に密集したメモリ集積回路において、隣接したアドレスを一定の範囲において飛び地を作りながら処理することを指します。

2.19.1. キャッシュライン

注記

詳しくは「表:ワークアイテム・グループ関数」(表B.102「表:ワークアイテム・グループ関数」)を参照ください。

メモリストライドは不連続にキャッシュの領域をまたぐことを指し、グローバルメモリ、コンスタントメモリにアクセスする際に発生します。メモリへのアクセスは例外なしに、連続的にアクセスするのが望ましいとされていますが、まさにストライド(キャッシュラインをまたぎ越す)が発生する状況です。

__global int* globalArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = globalArray[ gid ];

行ワークグループ: < 16, 1, 1 >を使うと、1 cachelineに16のワークアイテムを詰め込むことができます。これがint型の場合であれば、16*sizeof(int)、つまりL3キャッシュの領域64バイトを使用をします。これはメモリのフル帯域を使用するベストな例です。

図2.8 図:キャッシュライン(1 cacheline, 16 work-items)

width=600

__global int* globalArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = globalArray[ gid + 1 ];

この例では、メモリの読み込みはアラインされていません。そのため1つ余分なキャッシュラインを使用してしまいます。この例ではフル帯域の半分の読み込みパフォーマンスとなります。

図2.9 図:キャッシュライン(2 cacheline, 16 work-items)

width=600

__global int* globalArray = ...;
int tmp;
int gid_1 = get_global_id(0);
int gid_2 = get_global_id(1);
int gsize = get_global_size(0);
tmp = globalArray[ gid_1 + gid_2 * gsize ];

2次元ワークグループ(4, 2, 1)のケースでは、メモリの読み込みは下図のようになり、4つのキャッシュラインを使用します。この例ではフル帯域の1/4の読み込みパフォーマンスとなります。

図2.10 図:キャッシュライン(4 cacheline, 8 work-items)

width=100

__global int* globalArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = globalArray[ gid * 16 ];
//tmp = globalArray[ gid * 32 ];

このケースは各ワークアイテムが新たなcachelineを使用するメモリストライドが発生します。この例ではフル帯域の1/16の読み込みパフォーマンスとなります。

図2.11 図:キャッシュライン(4 cacheline, 4 work-items)

width=400

__global int* globalArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = globalArray[ gid * 8 ];

このケースは2つのワークアイテム新たなcachelineを使用するメモリストライドが発生します。この例ではフル帯域の1/8の読み込みパフォーマンスとなります。

図2.12 図:キャッシュライン(4 cacheline, 8 work-items)

width=400

2.19.2. バンク

注記

詳しくは「表:ワークアイテム・グループ関数」(表B.102「表:ワークアイテム・グループ関数」)を参照ください。

Bank Conflict(バンクコンフリクト)はローカルメモリ特有のアドレス空間におけるインデックス(バンク)の重複が、パフォーマンスを落とす現象です。

ローカルメモリはグローバルメモリのキャッシュ(LLC等)とは異なり、メモリの帯域は固有のバンクに割り振られています。グローバルメモリではキャッシュラインにアラインすることでフル帯域を使うことができるのに対して、バンクでは固有のIDがアサインされている限り、読み込み時にフル帯域を使用できます。

__local int* localArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = localArray[ gid ];

この例では16個の固有バンクを使います。バンクの重複がないためフル帯域を実現できます。

図2.13 図:フル帯域

width=600

__local int* localArray = ...;
int tmp;
int gid = get_global_id(0)
tmp = localArray[ gid + 1 ];

この例では16個の固有バンクを使いますので、フル帯域を実現できます。

図2.14 図:フル帯域

width=600

__local int* localArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = localArray[ gid & ~1 ];

この例では、8つの固有バンクですが、各バンクで同じアドレスを使い、バンクコンフリクトが発生しません。この例はやや特殊なのですが、フル帯域を実現できます。

図2.15 図:フル帯域

width=600

__local int* localArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = localArray[ gid * 2 ];

このケースでは、8つの固有なバンクを使用します。従ってフル帯域の半分の読み込みパフォーマンスとなります。

図2.16 図:半帯域

width=600

__local int* localArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = localArray[ gid * 16 ];

このケースはカラムにアクセスする際に頻繁におきます。ここでは、16の要素のストライドをすることで、バンクコンフリクトが発生します。このアクセスパターンでは読み込みでフル帯域の1/16のパフォーマンスとなります。

図2.17 図:1/16帯域(16 Bank Conflicts)

width=400

__local int* localArray = ...;
int tmp;
int gid = get_global_id(0);
tmp = localArray[ gid * 17 ];

このケースは前の例と同様にカラムにアクセスする際に頻繁におきます。ここでは、17の要素のストライドをすることで、バンクコンフリクトを避けています。このアクセスパターンでは読み込みでフル帯域を使用できます。

図2.18 図:フル帯域

width=400

Copyright 2018-2019, by Masaki Komatsu