データ並列プログラミングの詳細については「「データ並列プログラミングモデル」」、「「カーネルの実行」」、「「NDRange」」、「「ワークグループ」」、「「ワークアイテム関数」」を参照ください。
データ並列プログラミングについては、タスク並列プログラミングと異なり、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.java.
package com.book.jocl.data;
import static org.jocl.CL.*;
import java.io.File;
import java.net.URL;
import java.nio.file.Paths;
import java.util.Scanner;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_device_id;
import org.jocl.cl_event;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;
public class MultiplicationTest {
private static final int MAT_DIM = 2;
private static final int WORK_DIM = 2;
private static final String KERNEL_PATH = "mult.cl";
private static void print_error(String src_msg, int err)
{
final String[] err_msg = new String[]{
"CL_SUCCESS",
"CL_DEVICE_NOT_FOUND",
"CL_DEVICE_NOT_AVAILABLE",
"CL_COMPILER_NOT_AVAILABLE",
"CL_MEM_OBJECT_ALLOCATION_FAILURE",
"CL_OUT_OF_RESOURCES",
"CL_OUT_OF_HOST_MEMORY",
"CL_PROFILING_INFO_NOT_AVAILABLE",
"CL_MEM_COPY_OVERLAP",
"CL_IMAGE_FORMAT_MISMATCH",
"CL_IMAGE_FORMAT_NOT_SUPPORTED",
"CL_BUILD_PROGRAM_FAILURE",
"CL_MAP_FAILURE",
"CL_MISALIGNED_SUB_BUFFER_OFFSET",
"CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST",
"CL_COMPILE_PROGRAM_FAILURE ",
"CL_LINKER_NOT_AVAILABLE",
"CL_LINK_PROGRAM_FAILURE",
"CL_DEVICE_PARTITION_FAILED",
"CL_KERNEL_ARG_INFO_NOT_AVAILABLE",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"CL_INVALID_VALUE",
"CL_INVALID_DEVICE_TYPE",
"CL_INVALID_PLATFORM",
"CL_INVALID_DEVICE",
"CL_INVALID_CONTEXT",
"CL_INVALID_QUEUE_PROPERTIES",
"CL_INVALID_COMMAND_QUEUE",
"CL_INVALID_HOST_PTR",
"CL_INVALID_MEM_OBJECT",
"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
"CL_INVALID_IMAGE_SIZE",
"CL_INVALID_SAMPLER",
"CL_INVALID_BINARY",
"CL_INVALID_BUILD_OPTIONS",
"CL_INVALID_PROGRAM",
"CL_INVALID_PROGRAM_EXECUTABLE",
"CL_INVALID_KERNEL_NAME",
"CL_INVALID_KERNEL_DEFINITION",
"CL_INVALID_KERNEL",
"CL_INVALID_ARG_INDEX",
"CL_INVALID_ARG_VALUE",
"CL_INVALID_ARG_SIZE",
"CL_INVALID_KERNEL_ARGS",
"CL_INVALID_WORK_DIMENSION",
"CL_INVALID_WORK_GROUP_SIZE",
"CL_INVALID_WORK_ITEM_SIZE",
"CL_INVALID_GLOBAL_OFFSET",
"CL_INVALID_EVENT_WAIT_LIST",
"CL_INVALID_EVENT",
"CL_INVALID_OPERATION",
"CL_INVALID_GL_OBJECT",
"CL_INVALID_BUFFER_SIZE",
"CL_INVALID_MIP_LEVEL",
"CL_INVALID_GLOBAL_WORK_SIZE",
"CL_INVALID_PROPERTY",
"CL_INVALID_IMAGE_DESCRIPTOR",
"CL_INVALID_COMPILER_OPTIONS",
"CL_INVALID_LINKER_OPTIONS",
"CL_INVALID_DEVICE_PARTITION_COUNT",
};
int index = -err;
if (err != CL_SUCCESS) {
System.out.printf("Failed Message: %s - Error Code: %d\n", src_msg, err, err_msg[index]);
System.exit(-1);
}
}
/* サンプルデータの生成をします。 */
public static void construct_matrix(int m_dim, int n_dim, int p_dim, float[]A, float[]B, float[]C) {
for(int i = 0; i < m_dim; i++) {
for(int k = 0; k < p_dim; k++) {
A[i * p_dim + k] = (float)(i+1);
System.out.printf("%f\n", A[i * p_dim + k]);
}
}
for(int k = 0; k < p_dim; k++) {
for(int j = 0; j < n_dim; j++) {
B[k * n_dim + j] = (float)(k+1);
System.out.printf("%f\n", B[k * n_dim + j]);
}
}
}
/* コンテキストの生成をします。 */
public static void show_result(int m_dim, int n_dim, int p_dim, float[] C, double execution_time) {
for(int i = 0; i < m_dim*n_dim; i++) {
System.out.printf("C[%d] = %f\n",i,C[i]);
}
/* コンテキストの生成をします。 */
System.out.printf("Execution time in seconds: %f\n", execution_time / 1000000000);
}
public static void main(String[] args) throws Exception
{
/* 行列の宣言 */
float[] A; //(1)
float[] B; //(2)
float[] C; //(3)
/* 行列の各次元の宣言 */
int[] m_dim, n_dim, p_dim;
/* */
int err;
/* */
int a_size, b_size, c_size;
/* */
long[] global_size = new long[WORK_DIM]; //(4)
/* 次元の設定：ここではsquare（N*N）にします。 */
m_dim = new int[]{MAT_DIM}; //(5)
n_dim = new int[]{MAT_DIM}; //(6)
p_dim = new int[]{MAT_DIM}; //(7)
/* A in M by P : B in P by N : C in M by N dimensional spaces */
a_size = n_dim[0]*p_dim[0]; //(8)
b_size = p_dim[0]*m_dim[0]; //(9)
c_size = n_dim[0]*m_dim[0]; //(10)
/* OpenCL APIの変数 */
cl_device_id[] device = new cl_device_id[1];
cl_context context;
cl_command_queue queue;
cl_program program;
cl_platform_id[] platform = new cl_platform_id[1];
cl_kernel kernel;
cl_event profile_event = new cl_event();
/* グローバル空間の次元の宣言 */
int nd = WORK_DIM; //(11)
/* OpenCLバッファオブジェクトの宣言 */
cl_mem a_input; //(12)
cl_mem b_input; //(13)
cl_mem c_output; //(14)
/*
* OpenCLデバイスのプラットフォームの特定
* 最初に見つけたプラットフォームを使用します。
* */
clGetPlatformIDs(1, platform, null);
/*
* CPU/GPUデバイスの情報取得をします。
* */
err = clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_GPU, 1, device, null);
if(err == CL_DEVICE_NOT_FOUND) {
err = clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_CPU, 1, device, null);
}
if(err < 0) {
print_error("clGetDeviceIDs",err);
}
/* コンテキストの生成をします。 */
context = clCreateContext(null, 1, device, null, null, null);
/*
* OpenCL Cのソースコードをファイル（.cl）から
* 読み込みコンパイル・ビルドします
* */
/* ファイルを読み込みバッファーに投入します */
StringBuffer sb = new StringBuffer();
URL resource = MultiplicationTest.class.getResource(KERNEL_PATH);
String path = Paths.get(resource.toURI()).toFile().getAbsolutePath();
System.out.println(path);
Scanner sc = new Scanner(new File(path));
while(sc.hasNext()) {
sb.append(sc.nextLine() + "\n");
}
/* programの生成（ファイルからOpenCL Cソースコードの読み込み） */
program = clCreateProgramWithSource(context,
1, new String[]{sb.toString()}, null, null);
/* programのビルド（コンパイル）をします。 */
err = clBuildProgram(program, 0, null, null, null, null);
if(err < 0) {
print_error("clBuildProgram",err);
}
/* コマンドキューを生成します */
queue = clCreateCommandQueue(context, device[0], CL_QUEUE_PROFILING_ENABLE, null);
/* 行列のメモリー領域を割り当てます。 */
A = new float[a_size]; //(15)
B = new float[b_size]; //(16)
C = new float[c_size]; //(17)
/* 行列の初期化とサンプルデータの設定をします */
construct_matrix(m_dim[0], n_dim[0], p_dim[0], A, B, C); //(18)
/* OpenCLバッファーを生成します。
* a_inputとb_inputを入力、c_outputを出力とします。
* a_inputはA行列に対応します。
* */
a_input = clCreateBuffer(
context,
CL_MEM_READ_ONLY,
Sizeof.cl_float * a_size,
null, null); //(19)
b_input = clCreateBuffer(
context,
CL_MEM_READ_ONLY,
Sizeof.cl_float * b_size,
null, null); //(20)
c_output = clCreateBuffer(
context,
CL_MEM_WRITE_ONLY,
Sizeof.cl_float * c_size,
null, null); //(21)
/* Kernelを生成します */
kernel = clCreateKernel(program, "mult", null);
// Set the arguments to our compute kernel
/* */
err = 0;
err |= clSetKernelArg(kernel, 0, Sizeof.cl_int,Pointer.to(m_dim));
err |= clSetKernelArg(kernel, 1, Sizeof.cl_int,Pointer.to(n_dim));
err |= clSetKernelArg(kernel, 2, Sizeof.cl_int,Pointer.to(p_dim));
err |= clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(a_input));
err |= clSetKernelArg(kernel, 4, Sizeof.cl_mem, Pointer.to(b_input));
err |= clSetKernelArg(kernel, 5, Sizeof.cl_mem, Pointer.to(c_output));
if(err < 0) {
print_error("clSetKernelArg",err);
}
/* A行列とB行列をデバイスのバッファーにコピーします */
err = clEnqueueWriteBuffer(
queue,
a_input,
CL_TRUE,
0,
Sizeof.cl_float * a_size,
Pointer.to(A),
0,
null,
null); //(22)
if(err < 0) {
print_error("clEnqueueWriteBuffer on A",err);
}
clEnqueueWriteBuffer(
queue,
b_input,
CL_TRUE,
0,
Sizeof.cl_float * b_size,
Pointer.to(B),
0,
null,
null); //(23)
if(err < 0) {
print_error("clEnqueueReadBuffer on B",err);
}
/*
* 行列の次元をnとmとし、カーネルに回送する処理サイズパラメータに設定します
* 次元数が２となるので、nd変数は２とします。
* */
global_size[0] = (int) n_dim[0]; //(24)
global_size[1] = (int) m_dim[0]; //(25)
/*
* Kernelの実行をします。
* nd : グローバル空間の次元
* global_size : カーネルに回送する処理サイズ
* profile_event : 統計情報取得に使います
* */
err = clEnqueueNDRangeKernel(
queue,
kernel,
nd,
null,
global_size,
null,
0,
null,
profile_event); //(26)
if(err < 0) {
print_error("clEnqueueNDRangeKernel",err);
}
clFinish(queue);
/* 開始・終了処理時間の宣言をします。 */
long[] ev_start_time= new long[]{0};
long[] ev_end_time= new long[]{0};
double execution_time = 0.0;
/*
* 演算に要した処理時間をデバイスから取得します
* ev_start_time：カーネル演算開始時間
* ev_end_time：カーネル演算終了時間
* */
err = clGetEventProfilingInfo(profile_event,
CL_PROFILING_COMMAND_START,
Sizeof.cl_long,
Pointer.to(ev_start_time),
null); //(27)
err = clGetEventProfilingInfo(profile_event,
CL_PROFILING_COMMAND_END,
Sizeof.cl_long,
Pointer.to(ev_end_time),
null); //(28)
if(err < 0) {
print_error("clGetEventProfilingInfo",err);
}
/* デバイスから演算結果を読み込みます */
err = clEnqueueReadBuffer(
queue,
c_output,
CL_TRUE, 0,
Sizeof.cl_float * c_size,
Pointer.to(C), 0,
null,
null ); //(29)
if(err < 0) {
print_error("clEnqueueReadBuffer on C",err);
}
execution_time = ev_end_time[0] - ev_start_time[0];
/* 実行時間を表示します。 */
show_result(m_dim[0], n_dim[0], p_dim[0], C, execution_time);
/* OpenCL APIで割り当てたメモリーを解放します */
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseMemObject(a_input); //(30)
clReleaseMemObject(b_input); //(31)
clReleaseMemObject(c_output); //(32)
clReleaseCommandQueue(queue);
clReleaseContext(context);
}
}
行列Aを宣言。 | |
行列Bを宣言。 | |
行列Cを宣言。 | |
global_size（グローバル空間のサイズ）の次元を２次元に設定。 | |
Mを設定。 | |
Nを設定。 | |
Pを設定。 | |
Aの要素数を計算。 | |
Bの要素数を計算。 | |
Cの要素数を計算。 | |
グローバル空間の次元数を指定。 | |
Aのバッファオブジェクトを宣言。 | |
Bのバッファオブジェクトを宣言。 | |
Cのバッファオブジェクトを宣言。 | |
行列Aのメモリ領域を確保。 | |
行列Bのメモリ領域を確保。 | |
行列Cのメモリ領域を確保。 | |
行列Aと行列Bの初期化。 | |
Aのバッファオブジェクトを生成。 | |
Bのバッファオブジェクトを生成。 | |
Cのバッファオブジェクトを生成。 | |
初期化した行列Aをバッファオブジェクトにコピー。 | |
初期化した行列Bをバッファオブジェクトにコピー。 | |
Nを行空間の次元に設定。 | |
Mをカラム空間の次元に設定。 | |
clEnqueueNDRangeKernelを使いデバイスにカーネルを送る。profile_eventをカーネルと関連付ける。 | |
カーネル演算開始時間を取得。 | |
カーネル演算終了時間を取得。 | |
デバイスから演算結果を読み込み。 | |
Aのバッファオブジェクトのメモリ領域を解放。 | |
Bのバッファオブジェクトのメモリ領域を解放。 | |
Cのバッファオブジェクトのメモリ領域を解放。 |
mult.cl.
__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)
}
}
ホストプログラムで指定したm_dim変数 | |
ホストプログラムで指定したn_dim変数 | |
ホストプログラムで指定したp_dim変数 | |
Aのバッファオブジェクト | |
Bのバッファオブジェクト | |
Cのバッファオブジェクト | |
次元0のグローバルIDをiに代入 | |
次元1のグローバルIDをjに代入 | |
A(i,k)* B(k,j)を計算。 | |
C（i,j）を計算。 |
出力.
1.000000 1.000000 2.000000 2.000000 C[0] = 3.000000 C[1] = 3.000000 C[2] = 6.000000 C[3] = 6.000000 Execution time in seconds: 0.000011
/* A * B = C * * | 1 1 | * | 1 1 | = | 3 3 | * | 2 2 | | 2 2 | | 6 6 | * */
上記の出力は想定結果と一致することが確認できます。
この項目のプログラムではワークグループは一つの最もシンプルな設計となります。ワークグループを複数とする実装例を見るのだれば「ワークグループ実装例」（「実装例」）を参照ください。
「表：ランタイムエラー一覧」（表D.1「表：ランタイムエラー一覧」）と「表：コンパイルタイムエラー一覧」（表D.2「表：コンパイルタイムエラー一覧」）も参照ください。
JOCLでは以下のようにRuntimeExceptionを継承したCLExceptionが投げられます。
public class CLException extends java.lang.RuntimeException
例外を検知するにはこれを捕捉します。この例外は、カーネル関数でのエラーログも自動で集めてIDEなどに出力をします。
JOCLでは、コンパイルタイムエラーは原則としてJavaの例外処理に捉えられます。しかしkernel内のランタイムエラー（文法間違いではないがホストとの連携が不正等）などは実験したところ処理しませんでした。
このためエラーを検出するために以下のようなエラー表示関数を作っておくとよいでしょう。
private static void print_error(String src_msg, int err)
{
final String[] err_msg = new String[]{
"CL_SUCCESS",
"CL_DEVICE_NOT_FOUND",
"CL_DEVICE_NOT_AVAILABLE",
"CL_COMPILER_NOT_AVAILABLE",
"CL_MEM_OBJECT_ALLOCATION_FAILURE",
"CL_OUT_OF_RESOURCES",
"CL_OUT_OF_HOST_MEMORY",
"CL_PROFILING_INFO_NOT_AVAILABLE",
"CL_MEM_COPY_OVERLAP",
"CL_IMAGE_FORMAT_MISMATCH",
"CL_IMAGE_FORMAT_NOT_SUPPORTED",
"CL_BUILD_PROGRAM_FAILURE",
"CL_MAP_FAILURE",
"CL_MISALIGNED_SUB_BUFFER_OFFSET",
"CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST",
"CL_COMPILE_PROGRAM_FAILURE ",
"CL_LINKER_NOT_AVAILABLE",
"CL_LINK_PROGRAM_FAILURE",
"CL_DEVICE_PARTITION_FAILED",
"CL_KERNEL_ARG_INFO_NOT_AVAILABLE",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"",
"CL_INVALID_VALUE",
"CL_INVALID_DEVICE_TYPE",
"CL_INVALID_PLATFORM",
"CL_INVALID_DEVICE",
"CL_INVALID_CONTEXT",
"CL_INVALID_QUEUE_PROPERTIES",
"CL_INVALID_COMMAND_QUEUE",
"CL_INVALID_HOST_PTR",
"CL_INVALID_MEM_OBJECT",
"CL_INVALID_IMAGE_FORMAT_DESCRIPTOR",
"CL_INVALID_IMAGE_SIZE",
"CL_INVALID_SAMPLER",
"CL_INVALID_BINARY",
"CL_INVALID_BUILD_OPTIONS",
"CL_INVALID_PROGRAM",
"CL_INVALID_PROGRAM_EXECUTABLE",
"CL_INVALID_KERNEL_NAME",
"CL_INVALID_KERNEL_DEFINITION",
"CL_INVALID_KERNEL",
"CL_INVALID_ARG_INDEX",
"CL_INVALID_ARG_VALUE",
"CL_INVALID_ARG_SIZE",
"CL_INVALID_KERNEL_ARGS",
"CL_INVALID_WORK_DIMENSION",
"CL_INVALID_WORK_GROUP_SIZE",
"CL_INVALID_WORK_ITEM_SIZE",
"CL_INVALID_GLOBAL_OFFSET",
"CL_INVALID_EVENT_WAIT_LIST",
"CL_INVALID_EVENT",
"CL_INVALID_OPERATION",
"CL_INVALID_GL_OBJECT",
"CL_INVALID_BUFFER_SIZE",
"CL_INVALID_MIP_LEVEL",
"CL_INVALID_GLOBAL_WORK_SIZE",
"CL_INVALID_PROPERTY",
"CL_INVALID_IMAGE_DESCRIPTOR",
"CL_INVALID_COMPILER_OPTIONS",
"CL_INVALID_LINKER_OPTIONS",
"CL_INVALID_DEVICE_PARTITION_COUNT",
};
int index = -err;
if (err != CL_SUCCESS) {
System.out.printf("Failed Message: %s - Error Code: %d\n", src_msg, err, err_msg[index]);
System.exit(-1);
}
}本書ではコード行数が肥大するため、以降のサンプルコードではエラーコードの処理はしていません。
Copyright 2018-2019, by Masaki Komatsu