16.13. カーネル引数設定

Prev		Next

16.13.1. set_arg

set_arg関数はkernelの特定の引数に対して引数値を設定。

Note

詳しくは「表：clSetKernelArg」（Table B.81, “表：clSetKernelArg”）を参照ください。

set_arg(self,
    index, #(1)
    arg) #(2)

(1)	引数の添字（インデックス）を指定。カーネルが宣言している引数の数をnとした場合、左端の引数を0で始めてn - 1 までの添字でカーネルの引数の番号を指定。
(2)	indexで指定した引数に、カーネルの引数値として渡したいメモリオブジェクトまたはデータを指定。

引数の設計する際には、メモリ領域の種類（private, local, global）、どれだけのバイトサイズを割り当てるか、どのデータ型というように、カーネル関数側で決めうちしておかないと混乱します。

以下の項目では、最も初心者がはまる傾向があるメモリ領域の種類を見てみます。

プライベート引数

private修飾子を付けた引数はプライベート引数となります。最もシンプルなプライベート引数の宣言は以下のようにします。

__kernel void helloworld(uint p_num)
{

}

つまりプライベート引数は既定の設定のため、メモリ空間についての指定は不要となります。次に実装をしてみましょう。

import pyopencl as cl
import numpy as np

p_program = cl.Program(ctx, """
    __kernel void helloworld(__private uint p_num)
    {
            printf("My private number: %u\\n", p_num);
    }
    """).build()

これに対応したclSetKernelArg関数を考えてみましょう。ここではuintに対応するホストAPIの型として、cl_uintで変数を宣言します。

import pyopencl as cl
import numpy as np

num = np.int32(100) #(1)

p_kernel = cl.Kernel(p_program, name="helloworld")
p_kernel.set_arg(0, num)

cl.enqueue_task(queue, p_kernel)

(1)

numpy

出力.

My private number: 100

結果は100が表示されました。

local引数

ローカル引数を宣言する際はポインタにします。

__kernel void helloworld(__local uint* l_num)
{

}

ローカルメモリ空間内に割り当てる引数はホストAPIのLocalMemoryを使ってカーネル関数に渡します。

class pyopencl.LocalMemory(size)

sizeはローカルメモリで使う要素数を指定します。

kernel.set_arg(0, cl.LocalMemory(16))

このようにバッファオブジェクトを使わず明示的にローカルメモリとカーネルの引数に設定します。

例えばカーネル関数を以下のように設定した場合の挙動を見てみましょう。

import pyopencl as cl
import numpy as np

program = cl.Program(ctx, """
    __kernel void helloworld(__local uint* loc) {
        loc = 100;
            printf("My Local Number: %u\\n", loc);
    }
    """).build()

int型のローカルメモリ空間に100を設定して、それを出力します。

My Local Number: 100

カーネル関数内で初期化した通りの100が表示されます。

Warning

ローカル引数はホストAPIでNULLと設定します。ホストAPIからのデータ初期化ではなく、カーネル内で初期化をします。

global/constant引数

グローバル引数はグローバルメモリ空間（ヒープ領域に該当）に割り当てられるため、バッファオブジェクトを生成します。そのためカーネル関数内では常にポインタで宣言します。

g_program = cl.Program(ctx, """
    __kernel void helloworld(__global float* g_mat) {
        printf("Accessing mat[3]: %#f\\n", g_mat[3]);
    }
    """).build()

このカーネル関数に対応したホストAPIは以下のようにします。

ary = np.arange(16).astype(np.float32)

ary_mem = cl.Buffer(ctx, cl.mem_flags.USE_HOST_PTR, hostbuf=ary)

# g_programインスタンスの生成

g_kernel = cl.Kernel(g_program, name="helloworld")
g_kernel.set_arg(0, ary_mem)

cl.enqueue_task(queue, g_kernel)

グローバルメモリを使った引数を指定する際には、バッファオブジェクト（上記例では「g_mat」）を使います。

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

Accessing mat[3]: 3.000000

numpyのarangeで生成した値と同じものと確認できました。