タスク並列プログラミングの詳細については「the section called “タスク並列プログラミングモデル”」を参照ください。
この項目で以下のファイルを使います。
まず以下のようなカーネル関数をOpenCL-C言語で記述してみます。
カーネル関数.
__kernel void helloworld(__global int* number, char c) //(1) { int4 num4 = vload4(0, number + 1); //(2) float4 fl4 = (float4) (1.0f, 2.0f, 3.0f, 4.0f); //(3) printf("Hello World! Buddy!\n"); //(4) printf("%s\n", "Text literal is output just like this!"); printf("input integer is %d\n", number[4]); //(5) printf("15 in hex format is %#x\n", 15); printf("num4 vector is %#v4x\n", num4); //(6) printf("fl4 vector is %#v4x\n", fl4); printf("input character is %c\n", c); //(7) }
カーネルを宣言。第一引数にグローバル変数のnumber、第二引数にプライベート変数のbow、第三引数にプライベート変数のc。 | |
配列「number」のポインタをインクリメントした位置から4つのアドレスを、num4にコピー。「グローバル変数からプライベート変数へのロード」が行なわれます。 | |
printfはC言語とほぼ同じように使えます。 | |
「%s」は文字列の表示を指定します。「\n」は改行を指定します。 | |
「%d」は整数の表示を指定します。「\n」は改行を指定します。 | |
「%#v4x」は、v4つまり4つの要素からなるベクトルを、0xで始まる16進数型式での表示を指定します。num4にコピーした「number」は添字1から4つの要素ですが、「number」変数の添字5は未定義のため、0が表示されます。 | |
プライベートメモリ内の変数cにアクセスします。 |
次にこれに対応するスクリプトのコードをしてみます。
TaskTest.py.
import pyopencl as cl import numpy as np data = np.arange(16).astype(np.int32) devices = [cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0]] ctx = cl.Context(devices) queue = cl.CommandQueue(ctx) mf = cl.mem_flags a = np.arange(16).astype(np.int32) a_mem = cl.Buffer(ctx, cl.mem_flags.USE_HOST_PTR, hostbuf=a) c = np.uint8(90) program = cl.Program(ctx, """ __kernel void helloworld(__global int* number, char c) { int4 num4 = vload4(0, number + 1); float4 fl4 = (float4) (1.0f, 2.0f, 3.0f, 4.0f); printf("Hello World! Buddy!\\n"); printf("%s\\n", "Text literal is output just like this!"); printf("input integer is %x\\n", number[4]); printf("15 in hex format is %#x\\n", 15); printf("num4 vector is %#v4x\\n", num4); printf("fl4 vector is %#v4x\\n", fl4); printf("input character is %c\\n", c); } """).build() kernel = cl.Kernel(program, name="helloworld") kernel.set_arg(0, a_mem) kernel.set_arg(1, c) cl.enqueue_task(queue, kernel) queue.finish()
OpenCLでタスク並列プログラミングを指す場合は原則として、enqueue_task関数を使ったカーネルコマンドのキューへの挿入を意味し、以下のようにすればカーネルコマンドがデバイス内で発行し実行されます。
cl.enqueue_task(queue, kernel)
第1引数のqueueはコマンドキューオブジェクト、第2引数のkernelはカーネルオブジェクトです。
コマンドキューオブジェクトを生成するには、以下のようにコンテキストオブジェクト、デバイスオブジェクトを先行して生成しておく必要があります。
queue = cl.CommandQueue(ctx)
コンテキストオブジェクトを生成するには以下のようにclCreateContext関数の引数としてデバイスオブジェクトが必要となります。
ctx = cl.Context(devices)
デバイスオブジェクトの生成は、clGetDeviceIDs関数を使いますが、第一引数でプラットフォームオブジェクトが必要となります。
devices = [cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0]]
この一連の流れは「図:OpenCL API UMLクラス図」(Figure 5.1, “図:OpenCL API UMLクラス図”)にあるように、プラットフォームからデバイス、コンテキストからコマンドキューやバッファ、カーネルという依存性として理解できます。
サンプルプログラムによる出力は以下のようになります。
/Users/komatsu/JOCLBook/target/classes/com/book/jocl/task/basic_helloworld.cl Hello World! Buddy! Text literal is output just like this! input integer is 4 15 in hex format is 0xf num4 vector is 0x2,0x3,0x4,0 fl4 vector is 0x25ad905c,0x25ad9064,0x25ad906c,0x25ad9074 input character is p
Copyright 2018-2019, by Masaki Komatsu