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

__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)

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

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

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

__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 ];

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

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

__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)

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

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

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

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

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

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

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

図2.13 図：フル帯域

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

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

図2.14 図：フル帯域

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

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

図2.15 図：フル帯域

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

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

図2.16 図：半帯域

__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)

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

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

図2.18 図：フル帯域