CPU、FPGA、GPUといったヘテロジニアス・デバイスでプログラムを動作させるにはカーネルと呼ばれる関数をOpenCL C言語で記述する必要があります。このカーネルのコードは
といったように、clという拡張子のファイルに保管され、OpenCL C言語の文字列を読み込んでホストアプリケーションでビルドまたは、ビルド済みバイナリをロードして実行します。
OpenCL C言語は、C99(ISO/IEC 9899:1999)規格に準拠します。OpenCL-Cの特徴として、カーネルというホストアプリケーションから直接呼び出すことはできません。ホストアプリケーションはカーネルの実行をコマンドキューに挿入することで間接的な呼び出しが可能となります。
まず以下のOpenCL C言語で「Hello」と「Hello World!」という文字を出力するカーネル関数を見てみましょう。
kernel_helloworld.cl.
__kernel void helloworld(__global uchar16* bow) { printf("input character array elements are %c%c%c%c%c\n", (*bow).s0, (*bow).s1, (*bow).s2, (*bow).s3, (*bow).s4); printf("char2 vector is %v16c\n", *bow); }
このソースコードは例えば、kernel_helloworld.clというファイルに記述することができます。
注目頂きたい点としては、__kernelという修飾子が関数の導入箇所にあることです。全てのkernel関数はこの修飾子をつける必要があります。
また戻り値型がvoidとなっていますが、これもカーネル関数の特徴であり、常に戻り値型はvoidとします。最後に引数のbowですが、これはホストアプリケーションの中で、
byte[] message = new byte[] {'h','e','l','l','o',' ','w','o','r','l','d','!'}; cl_mem msg_mem = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, Sizeof.cl_uchar16, Pointer.to(message), null); clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(msg_mem));
というように文字配列の初期化を行なっています。uchar16というのは、OpenCL C言語の組込みデータ型であり、16の要素を持つ符号なし文字配列を指します。8ビットのため、Javaで対応するのはchar型ではなく、byte型となります。
この関数の出力は以下のようになります。
input character array elements are hello char16 vector is h,e,l,l,o, ,w,o,r,l,d,!,
char16という16の要素でなる文字ベクトルである、bowという変数を取り出して「Hello World!」という文字を出力できました。
ホストアプリケーションのソースコードは以下のようになります。
kernel_helloworld.java.
package com.book.jocl.kernel; 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; import com.book.jocl.task.TaskTest; public class KernelTest { private static final String KERNEL_PATH = "kernel_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 = KernelTest.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); byte[] message = new byte[] {'h','e','l','l','o',' ','w','o','r','l','d','!'}; cl_mem msg_mem = clCreateBuffer(context, CL_MEM_USE_HOST_PTR, Sizeof.cl_uchar16, Pointer.to(message), null); /* Kernelを生成します */ kernel = clCreateKernel(program, FUNC, null); err = 0; err |= clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(msg_mem)); /* * Kernelの実行をします。 * */ err = clEnqueueTask(queue, kernel, 0, null, null); if(err < 0) { print_error("clEnqueueTask",err); } clFinish(queue); /* OpenCL APIで割り当てたメモリーを解放します */ clReleaseProgram(program); clReleaseKernel(kernel); clReleaseCommandQueue(queue); clReleaseContext(context); } }
この例ではchar16というOpenCL-Cのベクトル型が登場しましたが、以降の項目で詳細の解説を行います。
Copyright 2018-2019, by Masaki Komatsu