14.3. HelloWorld(タスク並列プログラミング)

Note

タスク並列プログラミングの詳細については「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)
}

(1)

カーネルを宣言。第一引数にグローバル変数のnumber、第二引数にプライベート変数のbow、第三引数にプライベート変数のc。

(2)

配列「number」のポインタをインクリメントした位置から4つのアドレスを、num4にコピー。「グローバル変数からプライベート変数へのロード」が行なわれます。

(3)

printfはC言語とほぼ同じように使えます。

(4)

「%s」は文字列の表示を指定します。「\n」は改行を指定します。

(5)

「%d」は整数の表示を指定します。「\n」は改行を指定します。

(6)

「%#v4x」は、v4つまり4つの要素からなるベクトルを、0xで始まる16進数型式での表示を指定します。num4にコピーした「number」は添字1から4つの要素ですが、「number」変数の添字5は未定義のため、0が表示されます。

(7)

プライベートメモリ内の変数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クラス図”)にあるように、プラットフォームからデバイス、コンテキストからコマンドキューやバッファ、カーネルという依存性として理解できます。

14.3.1. 出力

サンプルプログラムによる出力は以下のようになります。

/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