詳しくは「表：clEnqueueNDRangeKernel」（表B.88「表：clEnqueueNDRangeKernel」）を参照ください。
int org.jocl.CL.clEnqueueNDRangeKernel(
cl_command_queue command_queue, //(1)
cl_kernel kernel, //(2)
int work_dim, //(3)
long[] global_work_offset, //(4)
long[] global_work_size, //(5)
long[] local_work_size, //(6)
int num_events_in_wait_list, //(7)
cl_event[] event_wait_list, //(8)
cl_event event) //(9)有効なコマンドキューを指定。 | |
有効なカーネルオブジェクトを指定。 | |
グローバルワークアイテム数とワークグループ内のワークアイテム数を決定する際の次元数を指定。 | |
要素数work_dimの配列でワークアイテムのグローバルIDを決定する際のオフセット値を指定。 | |
要素数work_dimの配列で、カーネル関数を実行するwork_dim次元のグローバルワークアイテムの個数を指定。 | |
要素数work_dimの配列で、カーネル関数を実行する各ワークグループを構成するワークアイテムの数（ワークグループのサイズとも呼ぶ）を指定。 | |
event_wait_listで指定したイベントオブジェクトの数を指定します。 | |
このコマンドが実行される前に完了されているイベントを指定 | |
カーネルコマンドを識別するイベントオブジェクトが戻され、コマンド完了の確認やコマンド完了の待機に使える。 |
clEnqueueNDRangeKernelを実装する際に一番判断に迷うのは、以下の３つの引数です。
オフセットは作業したいデータの位置や、アルゴリズムの設計により恣意的に開発者が決めるものですが、原則として(0,0,0)で支障はないかと思います。
グローバルワークサイズ引数についても同様ですが、一般にワークサイズは２の冪乗としてください。これはプロセッサの演算ユニットをなるべき多く使うための最も基本的な処方箋です。プロセッサを有効に使うためには、できるだけ多くのグローバルワークサイズ引数があったほうが良いです。
ローカルワークサイズ引数は、IntelのGPUであれば64〜128個とします。ローカルワークサイズを決定する際に覚えておきたい点は、ローカルワークサイズは、ワークグループ数を決めるということです。
例えば1024個のグローバルワークサイズに対して、128個のワークグループ数を考えてみましょう。この場合のワークグループ数は8個（1024/128）となります。しかしローカルワークサイズが大きいと、共有ローカルメモリのサイズが肥大化して、推奨されるメモリ使用量の閾値（Intelは4K）を超える可能性が高くなります。
詳細は 「ワークグループ数とワークグループサイズの最適化」（表2.13「ワークグループ数とワークグループサイズの最適化」） を参照ください。
work_helloworld.cxx.
package com.book.jocl.kernel_execution;
import static org.jocl.CL.*;
import java.io.File;
import java.net.URL;
import java.nio.ByteBuffer;
import java.nio.ByteOrder;
import java.nio.IntBuffer;
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;
import com.book.jocl.kernel.KernelTest;
public class KernelExecutionTest {
private static final String KERNEL_PATH = "work_helloworld.cl";
private static final String FUNC = "helloworld";
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];
cl_context context;
cl_command_queue queue;
cl_program program;
cl_platform_id[] platform = new cl_platform_id[1];
cl_kernel kernel;
/*
* OpenCLデバイスのプラットフォームの特定
* 最初に見つけたプラットフォームを使用します。
* */
err = 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);
}
/* コンテキストの生成をします。 */
context = clCreateContext(null, 1, device, null, null, null);
/*
* OpenCL Cのソースコードをファイル（.cl）から
* 読み込みコンパイルします
* */
/* ファイルを読み込みバッファーに投入します */
StringBuffer sb = new StringBuffer();
URL resource = KernelExecutionTest.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);
cl_mem resultObj = clCreateBuffer(
context,
CL_MEM_ALLOC_HOST_PTR,
Sizeof.cl_uint*2048,
null,
null
);
kernel = clCreateKernel(program, FUNC, null);
err = 0;
err |= clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(resultObj));
if(err < 0) {
print_error("clSetKernelArg",err);
}
int dim = 1;
long[] global_offset = new long[]{0,0,0};
long[] global_work_size = new long[]{2048,1,1};
long[] local_work_size = new long[]{64,1,1};
//cl_event event = new cl_event();
//cl_event[] events = new cl_event[] {event};
err = clEnqueueNDRangeKernel(
queue,
kernel,
dim,
global_offset,
global_work_size,
local_work_size,
0,
null,
null);
if(err < 0) {
print_error("clEnqueueNDRangeKernel",err);
}
//clWaitForEvents(0, events);
ByteBuffer result;
result = clEnqueueMapBuffer(
queue,
resultObj,
CL_TRUE,
CL_MAP_READ,
0,
Sizeof.cl_uint*2048,
0,
null,
null,
null);
clEnqueueUnmapMemObject(queue, resultObj, result, 0, null, null);
clFinish(queue);
result.order(ByteOrder.LITTLE_ENDIAN);
int sum = 0;
for(int i = 0; i < 2048; i++) {
int tmp = result.getInt();
sum+=tmp;
}
System.out.printf("%d\n", sum);
clReleaseMemObject(resultObj);
clReleaseProgram(program);
clReleaseKernel(kernel);
clReleaseCommandQueue(queue);
clReleaseContext(context);
}
}
work_helloworld.cl.
__kernel void helloworld(__global uint* globalObj)
{
size_t gid = get_global_id(0);
globalObj[gid] = 10;
}
出力.
20480
詳しくは「表：clEnqueueTask」（表B.89「表：clEnqueueTask」）、（タスク並列プログラミング）を参照ください。
int org.jocl.CL.clEnqueueTask(
cl_command_queue command_queue, //(1)
cl_kernel kernel, //(2)
int num_events_in_wait_list, //(3)
cl_event[] event_wait_list, //(4)
cl_event event) //(5)clEnqueueNativeKernel関数はOpenCLコンパイラを使い、コンパイルされてないネイティブC/C++関数を実行するコマンドを挿入します。
詳しくは「表：clEnqueueNativeKernel」（表B.90「表：clEnqueueNativeKernel」）を参照ください。
関数の定義は以下のようになります。
int org.jocl.CL.clEnqueueNativeKernel(
cl_command_queue command_queue, //(1)
EnqueueNativeKernelFunction user_func, //(2)
Object args, //(3)
long cb_args, //(4)
int num_mem_objects, //(5)
cl_mem[] mem_list, //(6)
Pointer[] args_mem_loc, //(7)
int num_events_in_wait_list, //(8)
cl_event[] event_wait_list, //(9)
cl_event event) //(10)有効なコマンドキューを指定。ネイティブユーザ関数は、CL_DEVICE_EXECUTION_CAPABILITIESをCL_EXEC_NATIVE_KERNEL機能に設定したデバイスで生成されるcommand-queue上で実行される。 | |
ホストが呼び出せる関数をを指定。 | |
user_funcを呼び出す際の引数リストを指すオブジェクトを指定。 | |
引数リストのバイトサイズを指定。argsとcb_argsバイトサイズが指すデータは複数されて、複製した領域はuser_funcに渡される。 | |
argsに渡されるバッファオブジェクトの個数を指定。 | |
num_mem_objects > 0の場合、有効なバッファオブジェクトのリストを指定 | |
argsが指すメモリオブジェクトのハンドルが格納される領域を指すポインタ。ユーザ関数が実行される前に、メモリオブジェクトのハンドルはグローバルメモリへのポインタと交換される。 | |
event_wait_list で指定したイベントオブジェクトの数を指定します。 | |
このコマンドが実行される前に完了されているイベントを指定 | |
カーネルコマンドを識別するイベントオブジェクトが戻され、コマンド完了の確認やコマンド完了の待機に使える。 |
Copyright 2018-2019, by Masaki Komatsu