CPU、FPGA、GPUといったヘテロジニアス・デバイスでプログラムを動作させるにはカーネルと呼ばれる関数をOpenCL C言語で記述する必要があります。このカーネルのコードはPythonのソースコード内で文字列として記述し、OpenCL C言語の文字列を読み込んでホストアプリケーションでビルドまたは、ビルド済みバイナリをロードして実行します。

OpenCL C言語は、C99(ISO/IEC 9899:1999)規格に準拠します。OpenCL-Cの特徴として、カーネルというホストアプリケーションから直接呼び出すことはできません。ホストアプリケーションはカーネルの実行をコマンドキューに挿入することで間接的な呼び出しが可能となります。

まず以下のOpenCL C言語で「01234」と「0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15」という数値を出力するカーネル関数を見てみましょう。

カーネル関数. 

__kernel void helloworld(__global uchar16* bow)
{
    printf("input array elements are %d%d%d%d%d\\n", (*bow).s0, (*bow).s1, (*bow).s2, (*bow).s3, (*bow).s4);
    printf("vector is %v16d\\n", *bow);
}

注目頂きたい点としては、__kernelという修飾子が関数の導入箇所にあることです。全てのkernel関数はこの修飾子をつける必要があります。

また戻り値型がvoidとなっていますが、これもカーネル関数の特徴であり、常に戻り値型はvoidとします。最後に引数のbowですが、これはホストアプリケーションの中で初期化したデータをデバイスの管理するメモリ空間にコピーしたものです。

import numpy as np
import pyopencl as cl

devices = [cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0]]

ctx = cl.Context(devices)
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags

bow = np.arange(16).astype(np.uint8)
bow_mem = cl.Buffer(ctx, mf.USE_HOST_PTR, hostbuf=bow)

program = cl.Program(ctx, """
    __kernel void helloworld(__global uchar16* bow)
    {
        printf("input array elements are %d%d%d%d%d\\n", (*bow).s0, (*bow).s1, (*bow).s2, (*bow).s3, (*bow).s4);
        printf("vector is %v16d\\n", *bow);
    }
    """).build()

kernel = cl.Kernel(program, name="helloworld")
kernel.set_arg(0, bow_mem)

cl.enqueue_task(queue, kernel)

このように配列の初期化をホストプログラム内で行なっています。

この関数の出力は以下のようになります。

input array elements are 01234
vector is 0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15

uchar16という16の要素でなる8bit整数を要素に持つベクトルである、bowという変数を取り出して「0,1,2,3,4,5,6,7,8,9,10,11,12,13,14,15」という配列を出力できました。

この例ではuchar16というOpenCL-Cのベクトル型が登場しましたが、以降の項目で詳細の解説を行います。