vloadn関数はアラインメントのとれてないベクトルのロードに使います。

gentypen vloadn(
    size_t offset,
    const global gentype *p)

`sizeof(gentypen)`バイトのデータ`（p + (offset * n)`のアドレスで読み込み）を返します。

gentypeがcharまたはucharのときは、`(p + (offset * n))`として計算されるアドレスは8-bitとして揃う(アラインメントをとる)必要があります。

gentypeがshortまたはushortのときは、16-bit長で揃えます。int、uint、floatのときは、32-bit長で揃えます。long、ulong、doubleのときは、64-bit長で揃えます。

vstoren関数はアラインメントのとれてないベクトルの書き込みに使います。

void vstoren(
    gentypen data,
    size_t offset,
    global gentype *p)

sizeof(gentypen)のデータを`(p + (offset * n))`のアドレスに書き込みます。

gentypeがcharまたはucharのときは、`(p + (offset * n))`として計算されるアドレスは8-bitとして揃う(アラインメントをとる)必要があります。

gentypeがshortまたはushortのときは、16-bit長で揃えます。int、uint、floatのときは、32-bit長で揃えます。long、ulong、doubleのときは、64-bit長で揃えます。

vload_halfnはアラインメントのとれてないベクトルのロードをし、floatのベクトルを返します。

floatn vload_halfn(
    size_t offset,
    const global half *p)

sizeof(halfn)バイトのデータ`（p + (offset * n)`のアドレスで読み込み）を返します。

読み込んだデータはhalfnとして解釈されます。そしてhalfnの値はfloatnに変換され、floatnの値が返されます。`（p + (offset * n)`のアドレスは、16-bit長で揃える必要があります。

カーネル内で同期点を持つには、barrier関数を使います。カーネルを実行する*ワークグループ内のワークアイテム*は、その他の全てのワークアイテムがバリアーを超えた処理（バリアー関数の後に記述された処理）に進むために、この関数を実行する必要があります。この関数はカーネルを実行する全てのワークアイテムによって呼び出されます。

以下が関数の定義です。

void barrier (
    cl_mem_fence_flags flags
)

flagsは以下２つの値を取ります。

メモリ操作の順序性はメモリーフェンスを使うため、mem_fence関数のメモリへのコミットの順序の属性を持ちます。

barrierはメモリの整合性と操作の同期を同時に達成します。ただしbarrierをループや分岐処理内で使うことは、ハング（デッドロックに類似した現状）によるクラッシュを発生させるので注意が必要です。ハングは以下の状況で発生します。

__kernel void helloworld(__global uint* g_mat,
                __local uint* l_mat)
{

        size_t gid = get_global_id(0);
        size_t lid = get_local_id(0);
        size_t group = get_group_id(0);
        size_t lsize = get_local_size(0);
        size_t wgsize= get_num_groups(0);

        __global uint* output;

        //printf("g_mat[%d] = %u\n", group * lsize + lid, *(g_mat+gid));

        uint sum = 0;

        for(int i = 1; i <= 1000; i++) {
                sum += group;
        }

        l_mat[lid] = sum;
        barrier(CLK_LOCAL_MEM_FENCE);

        //printf("l_mat[%u] = %u\n", lid, l_mat[lid]);

        if(lid == 0) {
                uint group_sum = 0;

                for(size_t i = 0; i < lsize; i++) {
                        group_sum += l_mat[i];
                }
                g_mat[group] = group_sum;
                //printf("output[%u] = %u\n", gid, g_mat[group]);
        }

        barrier(CLK_GLOBAL_MEM_FENCE);

        if(group == 0){

                uint output_sum = 0;

                for(size_t i = 0; i < wgsize; i++) {
                        output_sum += g_mat[i];
                }
                g_mat[wgsize] = output_sum;
                //printf("output_sum = %u\n", output_sum);
        }
}

*ワークアイテム内でメモリの整合性を取りたい場合*にはmem_fence関数を使い読み込み、書き込みのタイミング・順序を整列させます。

以下が関数の定義です。

void mem_fence (
    cl_mem_fence_flags flags
)

この関数はカーネルを実行するワークアイテムのロードや保存の前後の順序を整列させます。これはメモリーフェンスの後に記述されたロードや保存が行われる前に、ロードや保存はメモリーにコミットされることを意味します。

この他にも、読み込み・書き込みだけに限定したフェンスを設けることができる以下の関数があります。

void read_mem_fence (
    cl_mem_fence_flags flags)
void write_mem_fence (
    cl_mem_fence_flags flags)

flags引数についてはbarrier関数と同様に、CLK_LOCAL_MEM_FENCE CLK_GLOBAL_MEM_FENCEの２つの値を指定できます。

__kernel void helloworld(__global uint* g_mat1,
                __global uint* g_mat2)
{
    size_t id = get_global_id(0);
        g_mat2[a] = g_mat1[id];
        mem_fence(CLK_LOCAL_MEM_FENCE);
        g_mat2[b] = g_mat1[id+1];
}

明示的メモリフェンスとは、g_mat2のインデックスaのstore、g_mat1のインデックスidのloadの処理がメモリにコミットされた後に、g_mat2のインデックスbへのstoreとg_mat1のインデックスid+1のloadが処理されることを保証します。

非同期複製の関数には2つの機能があります。

これは以下の関数プロトタイプの引数をグローバル変数をローカル変数に変えるだけで使いわけられます。

event_t async_work_group_copy (
    __local gentype *dst, //(1)
    const __global gentype *src, //(2)
    size_t num_gentypes, //(3)
    event_t event) //(4)

__kernel void helloworld(__global uint* g_mat, __local uint* l_mat, uint num_rows, uint num_cols)
{

        size_t gid = get_global_id(0);
        size_t lid = get_local_id(0);
        size_t group = get_group_id(0);
        size_t lsize = get_local_size(0);

        printf("g_mat[%d] = %u\n", group * lsize + lid, *(g_mat+gid)); //(1)

        event_t my_evt;
        my_evt =  async_work_group_copy(
                        l_mat, //(2)
                        g_mat + (group * lsize), //(3)
                        4, //(4)
                        (event_t) 0
                );


        wait_group_events(1, &my_evt); //(5)

        g_mat[lsize * lid + group] = l_mat[lid]; //(6)

}

g_mat[8] = 8
g_mat[9] = 9
g_mat[10] = 10
g_mat[11] = 11
g_mat[12] = 12
g_mat[0] = 0
g_mat[13] = 13
g_mat[1] = 1
g_mat[14] = 14
g_mat[2] = 2
g_mat[15] = 15
g_mat[3] = 3
g_mat[4] = 4
g_mat[5] = 5
g_mat[6] = 6
g_mat[7] = 7
0 //(1)
4
8
12
1 //(2)
5
9
13
2 //(3)
6
10
14
3 //(4)
7
11
15

barrier関数はワークグループ内のスレッドの同期をとるため、*「ローカルな同期」*を行ないます。アトミック関数は*主に競合アクセスが発生するデータに対する「グローバルまたはローカル」なメモリ整合性*を確保するために使います。ごく一部の例外を除いて、int型、unsigned int型のみで提供されています。

おそらく一番分かりやすいアトミック関数は、atomic_addです。まずグローバルメモリへ加算をする種類を見てみましょう。

int atomic_add (
    volatile __global int *p,
    int val)

ポインタpが指すアドレスで32-bitの値（oldと呼称）を読みこみ、次に`(old + val)`を計算し、結果をpが指すアドレスに書き込みます。関数はoldを返します。

次に変数をアトミックにインクリメントする関数です。

int atomic_inc (
    volatile __global int *p)

この関数はポインタpが指すアドレスで32-bitの値（oldと呼称）を読みこみ、次に`(old + 1)`を計算し、結果をpが指すアドレスに書き込みます。関数はoldを返します。

次にローカルメモリの整合性をとる種類をみてみましょう。

int atomic_add (
    volatile __local int *p,
    int val)

基本的にはグローバルメモリのものと一緒なのですが、アトミック関数を使う場合には、ローカルメモリの初期化にbarrier関数を使うことが必要な場合があります。もちろんグローバルメモリでも順序性が必要な場合であればbarrierと組み合わせることができます。グローバル変数と同様に、類似したatomic_incも用意されています。

浮動小数点数でアトミックなスワップ操作を行なう関数もあります。

float atomic_xchg (
    volatile __global float *p,
    float val)

この関数はアドレスpに保存された古い値を、引数で指定した値valとスワップし、古い値を返します。この関数はint型でもサポートされています。

Figure 17.1. シャッフル

定義. 

gentypen shuffle (
    gentypem x,
    ugentypen mask)

shuffle関数は一つの入力ベクトルに対してマスクを適用し、要素間の順序変更をします。shuffleではマスク要素は`ilogb(2m-1)`数の最下位ビットのみを使います。

定義. 

gentypen shuffle2 (
    gentypem x,
    gentypem y,
    ugentypen mask)

shuffle2関数は２つの入力ベクトルに対してマスクを適用し、要素間の順序変更をします。shuffle2ではマスク要素は`ilogb(2m-1)+1`のみを使います。

__kernel void helloworld(
                __global uint* g_mat)
{

        float4 in1 = (float4)(1.5f, 15.0f,5.0f,7.5f);
        float4 in2 = (float4)(-20.0f, -5.0f, 1.0f, 100.0f);
        uint4 shuffle_mask = (uint4)(0,3,2,1);
        uint8 shuffle2_mask = (uint8)(0,3,2,1,5,4,7,6);

        float4 result1 = shuffle(in1, shuffle_mask);
        printf("Shuffle Result vector is %#2.1v4f\n", result1);

        float8 result2 = shuffle2(in1, in2, shuffle2_mask);
        printf("Shuffle2 Result vector is %#2.1v8f\n", result2);

}

Shuffle Result vector is 1.5,7.5,5.0,15.0
Shuffle2 Result vector is 1.5,7.5,5.0,15.0,-5.0,-20.0,100.0,1.0