データ並列プログラミングについては、タスク並列プログラミングと異なり、clEnqueueNDRangeKernel関数を使い、グローバル・ローカル空間の次元やサイズを指定する必要があります。

この実装例では「2 x 2」の行列の乗算をしてみます。式に起こすと、「A x B = C」とし、AとBを掛けてCを計算します。A、B、CはNxM(2 x 2)の行列とします。正確には、Aは(NxP)、Bは（PxM）、Cは(NxM)の実数空間とします。AとBは入力、Cは出力に使います。（つまりP=2、N=M=2）

行列はバッファオブジェクトとしてグローバルメモリ領域に保管します。このサンプルでは、ホストポインタを使った前の項目に対して、バッファの生成時は領域を空にして、バッファの書き込みコマンド（clEnqueueWriteBuffer関数）をキューに挿入します。また、出力行列Cについては、バッファの読み込みコマンド（clEnqueueReadBuffer関数）を使い演算結果を取得します。

またOpenCLのプロファイルフラグを有効とし、clGetEventProfilingInfo関数を使いカーネル実行時間を採集します。

MultiplicationTest.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
data_mem = cl.Buffer(ctx, mf.USE_HOST_PTR, hostbuf=data)

MAT_DIM = 2
WORK_DIM = 2 #(1)

# 次元の設定：ここではsquare（N*N）にします。

m_dim = MAT_DIM #(2)
n_dim = MAT_DIM #(3)
p_dim = MAT_DIM #(4)

# A in M by P : B in P by N : C in M by N dimensional spaces */

a_size = n_dim * p_dim #(5)
b_size = p_dim * m_dim #(6)
c_size = n_dim * m_dim #(7)

global_size = (WORK_DIM, WORK_DIM, ) #(8)

a = np.arange(a_size).astype(np.float32) #(9)
b = np.arange(b_size).astype(np.float32) #(10)
c = np.arange(c_size).astype(np.float32) #(11)

print("a: " + str(a))
print("b: " + str(b))

a_mem = cl.Buffer(ctx, cl.mem_flags.USE_HOST_PTR, hostbuf=a) #(12)
b_mem = cl.Buffer(ctx, cl.mem_flags.USE_HOST_PTR, hostbuf=b) #(13)
c_mem = cl.Buffer(ctx, cl.mem_flags.USE_HOST_PTR, hostbuf=c) #(14)

program = cl.Program(ctx, """
    __kernel void mult(
        const int m_dim,
        const int n_dim,
        const int p_dim,
        global float* A,
        global float* B,
        global float* C)
    {
        int k;
        int i = get_global_id(0);
        int j = get_global_id(1);
        float tmp;
        if((i < n_dim) && (j < m_dim)) {
            tmp = 0.0f;
            for(k=0; k<p_dim;k++) {
                tmp += A[i * n_dim + k] * B[k * p_dim + j];
            }
            C[i * n_dim + j] = tmp;
        }
    }
    """).build()

kernel = cl.Kernel(program, name="mult") #(15)

kernel.set_arg(0, np.int32(m_dim)) #(16)
kernel.set_arg(1, np.int32(n_dim)) #(17)
kernel.set_arg(2, np.int32(p_dim)) #(18)
kernel.set_arg(3, a_mem) #(19)
kernel.set_arg(4, b_mem) #(20)
kernel.set_arg(5, c_mem) #(21)

cl.enqueue_nd_range_kernel( #(22)
    queue, #(23)
    kernel, #(24)
    global_work_size=global_size, #(25)
    local_work_size=(1, 1, )) #(26)

out = np.ndarray(c_size).astype(np.float32) #(27)

cl.enqueue_read_buffer(queue, mem=c_mem, hostbuf=out) #(28)

print("OpenCL Output: " + str(out))

mat = np.array([[0.0, 1.0], [2.0, 3.0]])

out_check = np.dot(mat, mat)

print("Numpy Output: " + str(out_check))

カーネル関数. 

__kernel void mult(
        const int m_dim, //(1)
        const int n_dim, //(2)
        const int p_dim, //(3)
        __global float* A, //(4)
        __global float* B, //(5)
        __global float* C) //(6)
{
        int k;
        int i = get_global_id(0); //(7)
        int j = get_global_id(1); //(8)
        float tmp;
        if((i < n_dim) && (j < m_dim)) {
                tmp = 0.0f;
                for(k=0; k<p_dim;k++) {
                        tmp += A[i * n_dim + k] * B[k * p_dim + j]; //(9)
                }
        C[i * n_dim + j] = tmp; //(10)
        }
}

出力. 

a: [ 0.  1.  2.  3.]
b: [ 0.  1.  2.  3.]
OpenCL Output: [  2.   3.   6.  11.]
Numpy Output: [[  2.   3.]
 [  6.  11.]]

/* A * B = C
 *
 *   | 0   1 | * | 0   1 | = | 2  3  |
 *   | 2   3 |   | 2   3 |   | 6  11 |
 *
 */

上記の出力は想定結果と一致することが確認できます。

class pyopencl.Error

class pyopencl.MemoryError
class pyopencl.LogicError
class pyopencl.RuntimeError