{Python}GPGPUで行列積を求める

GPGPUで行列積を求めて行きたいと思います。PyOpenCLについてはあまり理解していないところが多いので、正確ではないと思います。
ただ、私はこう理解していて、こうすると動きます、という、メモです……、。

numpyで求める

行列積はnumpyのdotで求めることができます。
a=行列A
b=行列B
product = numpy.dot(a,b)

GPGPUで行列積を求める

GPGPUで演算を行うには、

  • Context
  • Queue
  • Buffer
  • Kernel
  • を記述することが必要になります。

    Context

    どのGPUを使用するかを指定します。
    引数にinteractive=Falseで自動選択。
    pyopencl.create_some_context(interactive=False)


    Queue

    Contextのなかに新しいqueueを作成します。
    pyopencl.CommandQueue(context)

    Buffer

    ホスト(CPU)とデバイス(GPU)間のデータ受け渡し用のBufferを作成します。

    カーネルに引数として渡す行列と、結果を受ける行列を作成

    ここでは、-9~9までのランダムの整数からなる4*4の行列を作成
    アウトプット行列は行列aと同じ大きさのから行列を作成

    size = 4
    a = numpy.random.randint(-9, 9, (size,size)).astype(numpy.int32)
    b = numpy.random.randint(-9, 9, (size,size)).astype(numpy.int32)
    dest = np.empty(a.shape, dtype=np.int32)

    bufferの作成

    用途によって、READ_ONLY, WRITE_ONLY, READ_WRITEを指定する。
    ここではインプットがREAD_ONLY,アウトプットbufferがWRITE_ONLYになっているが、
    行列のすべての要素になんらかの演算をする場合など、インプットデータを加工したものを戻り値として得たい場合はREAD_WRITEにする。
    hostbufを指定することでホスト->デバイスの転送が自動で行われる。

    a_buf = pyopencl.Buffer(context, mem_flags.READ_ONLY | mem_flags.COPY_HOST_PTR, hostbuf=a)
    b_buf = pyopencl.Buffer(context, mem_flags.READ_ONLY | mem_flags.COPY_HOST_PTR, hostbuf=b)
    dest_buf = pyopencl.Buffer(context, mem_flags.WRITE_ONLY, dest.nbytes)


    Kernel関数

    Kernel関数を記述します。Kernel関数内の演算はGPUで処理されます。
    Programクラスでmat_productというkernel name(関数名と同じようなもの)を付けて記述します。
    あとでこのkernel nameを使ってpythonプログラム内から呼び出します。

    get_global_id()

    このメソッドで取得できるのは現在カーネルで処理しているデータの全体での位置です。
    行列として考えると、
    get_global_id(0) (0 = dimension 0)が行番号
    get_global_id(1) (1 = dimension 1)が列番号
    にあたります。(として理解しています…)
    ※global_id、local_id、Work Groupなどについては、仕事の時に書いたノートが見つかり次第、また別のエントリで整理したいと思います……。

    queueは一次元なので、行番号と列番号で要素を指定することはできませんが、
    dest_index = j * n + i とすることで1次元中の要素番号を求めることができます。
    行列積の計算方法通りに求めた値を、アウトプット用の空行列に順に格納していくことで、aとbの行列積を生成します。

    build()

    最後で必ず、kernelのコードをビルドします。

    Kernel関数の呼び出し

    programを使って、bufferを引数として渡し、kernel nameを指定し実行します。
    ここでは、計算に行列のサイズを使っているので、一緒に引数として渡します。
    カーネルにスカラー値を渡すときはnumpy型である必要があるため、予め変換しておきます。

    n = numpy.int32(size)
    e = program.mat_product(queue, a.shape, None, a_buf, b_buf, dest_buf, n)

    カーネル関数の実行が全て終わるまで待ちます。

    e.wait()



    bufferの値をホストに移動

    wait()ですべてのカーネルの実行が終わったら、アウトプット用のbufferに入った結果をホストにコピーします。
    pyopencl.enqueue_copy(queue, dest, dest_buf)
    これでdestに行列積が格納されたことになります。
    以上です。

    ソースコード:

    出力値:

    matrix a
    [[ 0 -2  5 -7]
     [ 7 -4  2 -8]
     [-8 -3 -8  4]
     [-4 -5  2  3]]
    
    matrix b
    [[-7  8  5 -9]
     [-1 -9  8  4]
     [ 6  5 -1  7]
     [-8 -7 -9  8]]
    
    numpy_product result
    [[  88   92   42  -29]
     [  31  158   73 -129]
     [ -21 -105  -92   36]
     [  21    2  -89   54]]
    
    GPU_product result
    [[  88   92   42  -29]
     [  31  158   73 -129]
     [ -21 -105  -92   36]
     [  21    2  -89   54]]

    参考:Beginner’s Tutorial In PyOpenCL
    PyOpenCLハンズオン in kyoto.py

    Leave a Reply

    Your email address will not be published. Required fields are marked *