ホストプログラムはOpenCLデバイスでの処理実行のためのカーネルを送るコマンドを発行します。

コマンドを発行時に、ホスト内に指定された整数のインデックス空間を生成します。カーネルのインスタンスはこのインデックス空間の各ポイントを実行します。

カーネルの各インスタンスをワークアイテムとよび、各インスタンスは整数のインデックス空間の座標で識別できます。

この座標はワークアイテムのグローバルIDと呼びます。

ワークアイテムはスレッドと考えると分かりやすいかもしれません。カーネルコマンドが発行されるとワークアイテムのコレクションを生成します。

インデックス空間の整数（又はグローバルID）で識別される各ワークアイテムは、同一のカーネルが指定する命令シーケンスを実行します。

つまり同一のカーネルが、並列で処理され、実行中の各カーネルはグローバルIDで識別することが可能となります。

各ワークアイテムの処理内容は多くの場合に異なります。

グローバルIDを入力データ配列要素のインデックスとして使用し、同じ命令シーケンスでも他のワークアイテムが持つグローバルIDが指す配列要素とはインデックスが異なる要素を処理したり、命令シーケンス内の分岐処理などによって、処理結果が各ワークアイテムで異なる形態をとらせることができます。

この処理方法をSPMD（Single Program Multiple Data）と呼ぶ事があります。

ワークアイテムはワークグループに分類できます。

ワークグループは整数空間を分割して、グローバルIDという単一の整数空間でなく、ワークグループIDとワークグループ内要素を識別するローカルIDの組み合わせによりワークアイテムの座標を表すことができます。

ワークグループはOpenCLでの並列処理の核心をなす概念です。

これは単にグローバル空間を削減するという効果ではなく、プラットフォームモデルでご案内したデバイスの構造をフルに活用するために必須なものです。

ワークグループ内のワークアイテムは、CU(Compute Unit)内のPE(Processing Element)で並行または並列に実行することができるのです。

ただしベンダーが提供するOpenCL実装ライブラリが並行・並列に処理するかは保証されていません。

カーネルインスタンスのインデックス空間は、N次元の範囲（Range）として定義できます。

ワークアイテムのグローバルIDとローカルIDは、N-タプルとして表すことができます。Nは1,2,3のいずれかを設定することができます。

例えば N=1 の場合は (0)、(1)、(2)、．．． といったようにインデックスは単なる整数です。

N=2 の場合は (0,0)、(0,1)、(0,2)、．．． というように2次元のインデックスで表せます。

N=3 の場合は (0,0,0)、(0,0,1)、(0,0,2) といったように3次元のインデックスで表せます。

ワークグループについても、同様にN-タプルで表現できます。例えば2次元のグローバルIDが(3,3)だとするとワークグループの(1,1)のローカルIDは(1,1)となります。この例は後ほど表を使って解説します。

それでは1次元のカーネルインスタンスの場合の例を見てみましょう。以下の表では、ワークアイテムの総数を8個とし、ワークグループを2つにした場合の、グローバルIDとローカルIDを列挙しています。

ローカルIDは(0,1,2,3,0,1,2,3)となり各ワークグループの固有なIDが割り振られます。

次に2次元の例を再度見てみましょう。2次元のインデックス空間を使う場合は大抵は2次元画像の処理をすることを踏まえて、座標を(x,y)の2-タプルで表してみましょう。

以下は解説で使う変数の定義です。

ローカルアイテム数については以下のように、ワークアイテム数をワークグループ数で割る事をで算出できます。

グローバルIDは、ワークグループIDとローカルアイテム数とローカルIDを合成して算出できます。

ワークグループIDと、ローカルIDの計算は以下のような数式を使います。

ワークグループIDの計算式は規格に沿ったものですが、少数部分を切り下げる床関数を使って整数に変換する必要があります。

最後に以下の表が、4x4の2次元インデックス空間の場合の各変数の推移となります。

OpenCL実行モデルは、カーネルインスタンスの識別IDやインデックス空間の定義の他にも、OpenCLのホストプログラムのハブとなる（アクセスパタンを指定する）コンテキストを定義します。

OpenCLアプリケーションでは、コンテキストは以下の４つのリソースを使ってカーネルの実行環境を定義することに用います。

DevicesとKernelsについては既に解説してますので、Programオブジェクトとメモリオブジェクトを見てみましょう。

Programオブジェクトはホストプログラム内でランタイムにビルド（カーネルソースコードのコンパイル）されます。

これはOpenCLデバイスのターゲットが分からない状態を前提とすれば、プログラムの実行時にコンパイルして各種デバイスに対して柔軟な運用ができるメリットがあります。

Programオブジェクトのビルドをするには、ソースが必要ですが、文字列型にしてホストプログラムに置いておくか、外部ファイルをロードしてきます。

最後にコンテキスト内においてメモリを生成・初期化することができ、生成されたメモリをメモリオブジェクトと呼びます。

メモリオブジェクトにはバッファオブジェクトとイメージオブジェクトの２つがあります。

ホストからデバイスへのやり取りは、コマンドキュー(command-queues)というキューに、処理したいコマンドを挿入して行ないます。

一つ目のカーネル実行コマンドはOpenCLデバイスのPE上でカーネルを実行させます。

二つ目は、コマンドキュー(command-queues)を介してデバイスにデータを転送したり、デバイスからデータを受け取る、データをマップ、アンマップするコマンドとなります。

データの型式はメモリオブジェクトというOpenCLがサポートするオブジェクトに変換します。

三つ目の同期コマンドは、コマンドが処理する順序を制御するために使います。

Contextは複数のコマンドキューと関連付けることができ、各コマンドキューは独立して機能し明示的に同期をさせるメカニズムは存在しません。

キューは順序通りコマンドを実行するFIFOとできます。反面、out-of-order実行（順不同実行）を行なえるコマンドを挿入することも可能です。

順不同実行では、コマンドは順序通り挿入され実行されていきますが、前のコマンドの終了を待機せずに実行をしていきます。

同一のコマンドキューでは、キューに挿入されたコマンドがイベントオブジェクトを生成することで、コマンド間の同期を行なうことができます。

同一カーネルが並列に処理されるが各カーネルインスタンスは固有の命令カウンタを持つことによって、例えば２次元のイメージ画像のピクセルをx軸、y軸として識別し、各ピクセルに一つのスレッド（ワークアイテム）を割り当てます。

SPMD（Single Program Multiple Data）はカーネルがデータと所有するプログラムカウンターに対して複数の処理要素（プロセッシングエレメント、PE）を使い並行で実行されるプログラミングをさします。つまり全ての演算リソースは同一カーネルを実行しますが、自らの命令カウンタを保ちます。つまり各カーネルのインスタンス（ワークアイテム）は互いに独立しています。

データ並列プログラミングモデルを実行するには、処理をさせたいカーネルをキューに挿入する必要があります。それには以下の関数を使います。

int org.jocl.CL.clEnqueueNDRangeKernel(
    cl_command_queue command_queue, //(1)
    cl_kernel kernel, //(2)
    int work_dim, //(3)
    long[] global_work_offset,
    long[] global_work_size, //(4)
    long[] local_work_size, //(5)
    int num_events_in_wait_list,
    cl_event[] event_wait_list,
    cl_event event)

キューに挿入後にカーネルは実行されます。カーネルはOpenCL-C言語で定義・記述する関数ですが、C言語との差異はあまりありません。clEnqueueNDRangeKernel関数はカーネルとホストプログラムの橋渡しをするだけではなく、データ並列プログラミングモデルをカーネルで実装するための不可欠な関数です。

タスク並列プログラミングは、C言語、C++で使うマルチスレッドと基本的には同じ意味を持ちます。

OpenCLではタスク並列の「タスク」を一つだけのワークアイテムを実行するカーネルと定義しています。

タスクを実行するには、処理をさせたいカーネルをキューに挿入する必要があります。それには以下の関数を使います。

int org.jocl.CL.clEnqueueTask(
    cl_command_queue command_queue, //(1)
    cl_kernel kernel, //(2)
    int num_events_in_wait_list, //(3)
    cl_event[] event_wait_list,
    cl_event event)

タスク並列の「並列」については以下のような可能性があります。

いずれの方式でタスク並列プログラミングを実行するかは設計上の問題とはなりますが、NDRangeの次元1、ワークアイテム1というシンプルなタスクのプログラミングモデルなので、用途はデータ並列プログラミングモデルと比べると柔軟に選べます。