タスク並列プログラミングの詳細については「「タスク並列プログラミングモデル」」を参照ください。
前の項目から引き継ぎ以下の３つのファイルを使います。
「.cl」拡張子はカーネルのソースコードを指します。
basic_helloworld.cl（カーネルです。）.
__kernel void helloworld(__global int* number, char c) //(1) { int4 num4 = vload4(0, number + 1); //(2) float4 fl4 = (float4) (1.0f, 2.0f, 3.0f, 4.0f); //(3) printf("Hello World! Buddy!\n"); //(4) printf("%s\n", "Text literal is output just like this!"); printf("input integer is %d\n", number[0]); //(5) printf("15 in hex format is %#x\n", 15); printf("num4 vector is %#v4x\n", num4); //(6) printf("fl4 vector is %#v4x\n", fl4); printf("input character is %c\n", c); //(7) }
カーネルを宣言。第一引数にグローバル変数のnumber、第二引数にプライベート変数のbow、第三引数にプライベート変数のc。 | |
配列「number」のポインタをインクリメントした位置から４つのアドレスを、num4にコピー。「グローバル変数からプライベート変数へのロード」が行なわれます。 | |
printfはC言語とほぼ同じように使えます。 | |
「%s」は文字列の表示を指定します。「\n」は改行を指定します。 | |
「%d」は整数の表示を指定します。「\n」は改行を指定します。 | |
「%#v4x」は、v4つまり４つの要素からなるベクトルを、0xで始まる16進数型式での表示を指定します。num4にコピーした「number」は添字1から４つの要素ですが、「number」変数の添字５は未定義のため、0が表示されます。 | |
プライベートメモリ内の変数cにアクセスします。 |
TaskTest.java.
package com.book.jocl.task;
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_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;
public class TaskTest {
private static final String KERNEL_PATH = "basic_helloworld.cl"; //(1)
private static final String FUNC = "helloworld"; //(2)
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 main(String[] args) throws Exception
{
int err;
/* OpenCL APIの変数 */
cl_device_id[] device = new cl_device_id[1]; //(3)
cl_context context; //(4)
cl_command_queue queue; //(5)
cl_program program; //(6)
cl_platform_id[] platform = new cl_platform_id[1]; //(7)
cl_kernel kernel; //(8)
/*
* OpenCLデバイスのプラットフォームの特定
* 最初に見つけたプラットフォームを使用します。
* */
err = clGetPlatformIDs(1, platform, null); //(9)
/*
* CPU/GPUデバイスの情報取得をします。
* */
err = clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_GPU, 1, device, null); //(10)
if(err == CL_DEVICE_NOT_FOUND) {
err = clGetDeviceIDs(platform[0], CL_DEVICE_TYPE_CPU, 1, device, null);
}
/* コンテキストの生成をします。 */
context = clCreateContext(null, 1, device, null, null, null); //(11)
/*
* OpenCL Cのソースコードをファイル（.cl）から
* 読み込みコンパイルします
* */
/* ファイルを読み込みバッファーに投入します */
StringBuffer sb = new StringBuffer();
URL resource = TaskTest.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ソースコードの読み込み） */
int[] errPtr = new int[1];
program = clCreateProgramWithSource(context,
1, new String[]{sb.toString()}, null, errPtr);
if(errPtr[0] < 0) {
print_error("clCreateProgramWithSource",errPtr[0]);
}
/* programのビルド（コンパイル）をします。 */
err = clBuildProgram(program, 0, null, null, null, null);
if(err < 0) {
print_error("clBuildProgram",err);
}
/* コマンドキューを生成します */
queue = clCreateCommandQueue(context, device[0], 0, null); //(12)
int[] myNumber = new int[4]; //(13)
myNumber[0] = 1;
myNumber[1] = 2;
myNumber[2] = 3;
myNumber[3] = 4;
char[] c = new char[]{'p'}; //(14)
cl_mem num_input = clCreateBuffer(context,
CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, //(15)
Sizeof.cl_int*4, //(16)
Pointer.to(myNumber), //(17)
null);
/* Kernelを生成します */
kernel = clCreateKernel(program, FUNC, null); //(18)
err = 0;
err |= clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(num_input)); //(19)
err |= clSetKernelArg(kernel, 1, Sizeof.cl_char, Pointer.to(c)); //(20)
if(err < 0){
print_error("clSetKernelArg",err);
}
/*
* Kernelの実行をします。
* */
err = clEnqueueTask(queue, kernel, 0, null, null); //(21)
if(err < 0) {
print_error("clEnqueueTask",err);
}
clFinish(queue); //(22)
/* OpenCL APIで割り当てたメモリーを解放します */
clReleaseProgram(program); //(23)
clReleaseKernel(kernel); //(24)
clReleaseCommandQueue(queue); //(25)
clReleaseContext(context); //(26)
clReleaseMemObject(num_input); //(27)
}
}
カーネルのソースファイル名を設定。カーネルとホストコードのソースは同一フォルダにある前提。 | |
カーネル関数の名称を設定。 | |
デバイスオブジェクトを宣言。 | |
コンテキストオブジェクトの宣言。 | |
コマンドキューオブジェクトの宣言。 | |
プログラムオブジェクトの宣言。 | |
プラットフォームオブジェクトの宣言。 | |
カーネルオブジェクトの宣言。 | |
OpenCLデバイスのプラットフォームの特定。 | |
CL_DEVICE_TYPE_GPUを指定して、GPUデバイスを検索します。取得したデバイスオブジェクトはdeviceオブジェクトを指すポインタに設定されます。 | |
コンテキストの生成をします。deviceオブジェクトと関連付けます。 | |
コマンドキューを生成します。 | |
整数型配列「myNumber」のメモリ空間内の領域を確保します。 | |
文字型の変数cに'p'を代入します。 | |
clCreateBufferで（カーネルのグローバル引数として使うため）バッファオブジェクトを生成。「CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR」フラグはメモリを読み込み専用、ホスト内のメモリ領域にあるポインタをデバイスに複製することを指定します。 | |
バッファオブジェクトに割り当てるバイトサイズを指定します。 | |
ホストポインタを指定します。 | |
カーネルオブジェクトを生成します。このカーネルは「FUNC」マクロで指定した「helloworld」関数を呼ぶことができます。 | |
グローバルメモリを使ってカーネルに引数を渡します。（グローバル領域に割り当てたカーネル引数は常にバッファオブジェクトとして渡します。） | |
プライベートメモリを使ってカーネルに引数を渡します。 | |
コマンドキューにカーネルコマンドを挿入します。第２引数で指定したカーネルが、コマンドキューに関連付けられたデバイス上で実行されます。 | |
先行してコマンドキューに挿入したOpenCLコマンドが終了するまでブロックします。 | |
プログラムオブジェクトのメモリ領域を解放します。 | |
カーネルオブジェクトのメモリ領域を解放します。 | |
コマンドキューオブジェクトのメモリ領域を解放します。 | |
コンテキストオブジェクトのメモリ領域を解放します。 | |
バッファオブジェクトのメモリ領域を解放します。 |
OpenCLでタスク並列プログラミングを指す場合は原則として、clEnqueueTask関数を使ったカーネルコマンドのキューへの挿入を意味し、以下のようにすればカーネルコマンドがデバイス内で発行し実行されます。
err = clEnqueueTask(queue, kernel, 0, null, null);
第１引数のqueueはコマンドキューオブジェクト、第2引数のkernelはカーネルオブジェクトです。
コマンドキューオブジェクトを生成するには、以下のようにコンテキストオブジェクト、デバイスオブジェクトを先行して生成しておく必要があります。
queue = clCreateCommandQueue(
context,
device[0],
0,
null);コンテキストオブジェクトを生成するには以下のようにclCreateContext関数の引数としてデバイスオブジェクトが必要となります。
context = clCreateContext(
null,
1,
device,
null,
null,
null);デバイスオブジェクトの生成は、clGetDeviceIDs関数を使いますが、第一引数でプラットフォームオブジェクトが必要となります。
err = clGetDeviceIDs(
platform[0],
CL_DEVICE_TYPE_GPU,
1,
device, null);この一連の流れは「図：OpenCL API UMLクラス図」（図2.1「図：OpenCL API UMLクラス図」）にあるように、プラットフォームからデバイス、コンテキストからコマンドキューやバッファ、カーネルという依存性として理解できます。
サンプルプログラムによる出力は以下のようになります。
/Users/komatsu/JOCLBook/target/classes/com/book/jocl/task/basic_helloworld.cl Hello World! Buddy! Text literal is output just like this! input integer is 1 15 in hex format is 0xf num4 vector is 0x2,0x3,0x4,0 fl4 vector is 0x25ad905c,0x25ad9064,0x25ad906c,0x25ad9074 input character is p
Copyright 2018-2019, by Masaki Komatsu