2.3.2. メモリー・アクセス・パターンの例の変更
以下は、単純なOpenCLカーネルのコード例です。
kernel void big_lmem_4r_4w_nosplit (global int* restrict in,
global int* restrict out) {
local int lmem[4][1024];
int gi = get_global_id(0);
int gs = get_global_size(0);
int li = get_local_id(0);
int ls = get_local_size(0);
int res = in[gi];
#pragma unroll
for (int i = 0; i < 4; i++) {
lmem[i][(li*i) % ls] = res;
res >>= 1; }
// Global memory barrier
barrier(CLK_GLOBAL_MEM_FENCE);
res = 0;
#pragma unroll
for (int i = 0; i < 4; i++) {
res ^= lmem[i][((ls-li)*i) % ls]; }
out[gi] = res;
}
この例のシステム・ビューア・レポートは、ストール可能なロードおよびストアを強調表示します。
図 16. 例のシステムビュー
図 17. 実施例のエリア・リポート
図 18. 例のカーネル・メモリー・ビューア
ロードとストアの動作の間の最初のバンクでは、アービトレーションが高い2つのメモリーバンクしか作成されないことに注意してください。次に、次のコード例に示すように、バンク・インデックスを第2次元に切り替えます。
kernel void big_lmem_4r_4w_nosplit (global int* restrict in,
global int* restrict out) {
local int lmem[1024][4];
int gi = get_global_id(0);
int gs = get_global_size(0);
int li = get_local_id(0);
int ls = get_local_size(0);
int res = in[gi];
#pragma unroll
for (int i = 0; i < 4; i++) {
lmem[(li*i) % ls][i] = res;
res >>= 1;
}
// Global memory barrier
barrier(CLK_GLOBAL_MEM_FENCE);
res = 0;
#pragma unroll
for (int i = 0; i < 4; i++) {
res ^= lmem[((ls-li)*i) % ls][i];
}
out[gi] = res;
}
カーネル・メモリー・ビューアでは、4つのメモリーバンクが別々のロード・ストア・ユニットで作成されていることがわかります。すべてのロードストア命令はストールフリーです。
図 19. バンク・インデックスを変更した後の例のカーネル・メモリー・ビューア
図 20. バンク・インデックスを変更した後の例のエリア・レポート