OpenCL-C組込み関数はstdio.hやstdlib.h、math.hといった標準ライブラリをインクルードできないので、それらの代用という面もありますが、ベクトルや並列演算、FPU等のハードウェアに合わせた機能を提供します。
詳しくは「表：ワークアイテム・グループ関数」（表B.102「表：ワークアイテム・グループ関数」）を参照ください。
get_global_id関数はdimindexで指定した次元で固有のグローバル・ワークアイテムIDを返します。
定義.
size_t get_global_id(uint dimindx)
サンプル.
size_t global_id_0 = get_global_id(0); size_t global_id_1 = get_global_id(1); size_t global_id_2 = get_global_id(2);
get_local_id関数はdimindexで指定した次元で固有のローカル・ワークアイテムIDを返します。
size_t get_global_offset(uint dimindx)
dimindexで指定した次元のオフセット（offset）を返します。clEnqueueNDRangeKernelのglobal_work_offset引数で指定したものが該当します。
size_t get_local_id(uint dimindx)
dimindexで指定した次元で固有のローカル・ワークアイテム（work item）IDを返します。
size_t get_group_id(uint dimindx)
dimindexで指定した次元のオフセット（offset）を返します。clEnqueueNDRangeKernelのglobal_work_offset引数で指定したものが該当します。
size_t get_global_size(uint dimindx)
dimindexで指定した次元のローカル・ワークアイテム（work item）の数を返します。clEnqueueNDRangeKernelのlocal_work_size引数で指定したものが該当します。
size_t get_local_size(uint dimindx)
dimindexで指定した次元のローカル・ワークアイテム（work item）の数を返します。clEnqueueNDRangeKernelのlocal_work_size引数で指定したものが該当します。
size_t get_num_groups(uint dimindx)
dimindexで指定した次元のワークグループ（work group）の数を返します。clEnqueueNDRangeKernelのlocal_work_size引数で指定したものが該当します。
__kernel void ndrange_id(__global write_only float *output) {
size_t global_id_0 = get_global_id(0); //(1)
size_t global_id_1 = get_global_id(1); //(2)
size_t global_id_2 = get_global_id(2); //(3)
size_t global_size_0 = get_global_size(0); //(4)
size_t global_size_1 = get_global_size(1); //(5)
size_t offset_0 = get_global_offset(0); //(6)
size_t offset_1 = get_global_offset(1); //(7)
size_t offset_2 = get_global_offset(2); //(8)
size_t local_id_0 = get_local_id(0); //(9)
size_t local_id_1 = get_local_id(1); //(10)
size_t local_id_2 = get_local_id(2); //(11)
int index_0 = global_id_0 - offset_0; //(12)
int index_1 = global_id_1 - offset_1; //(13)
int index_2 = global_id_2 - offset_2; //(14)
int index = index_2 * global_size_0 * global_size_1 + index_1 * global_size_0 + index_0;
float f = global_id_0 * 100.0f + global_id_1 * 10.0f + global_id_2 * 1.0f;
f += local_id_0 * 0.1f + local_id_1 * 0.01f + local_id_2 * 0.001f;
output[index] = f; //(15)
}次元1のグローバルIDを取得 | |
次元2のグローバルIDを取得 | |
次元3のグローバルIDを取得 | |
次元1のグローバルサイズ（ワークアイテム数）を取得 | |
次元2のグローバルサイズ（ワークアイテム数）を取得 | |
次元1のグローバルオフセットを取得 | |
次元2のグローバルオフセットを取得 | |
次元3のグローバルオフセットを取得 | |
次元1のローカルIDを取得 | |
次元2のローカルIDを取得 | |
次元3のローカルIDを取得 | |
グローバルIDとオフセットを調整 | |
グローバルIDとオフセットを調整 | |
グローバルIDとオフセットを調整 | |
出力変数「output」に代入 |
ホストプログラム側の設定は、64個のワークアイテムを3次元に分割するものとします。
package com.book.jocl.workitem;
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 WorkItemTest {
private static final String KERNEL_PATH = "ndrange_id.cl";
private static final String KERNEL = "ndrange_id";
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 = WorkItemTest.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);
float[] result = new float[64];
cl_mem res_buff = clCreateBuffer(
context,
CL_MEM_WRITE_ONLY | CL_MEM_USE_HOST_PTR,
Sizeof.cl_float*64, Pointer.to(result), null);
if(err < 0) {
print_error("clCreateBuffer", err);
}
kernel = clCreateKernel(program, KERNEL, null);
clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(res_buff));
int dim = 3;
long[] global_offset = new long[]{0, 0, 0};
long[] global_size = new long[]{4, 4, 4};
long[] local_size = new long[]{2, 2, 2};
err = clEnqueueNDRangeKernel(queue, kernel, dim, global_offset,
global_size, local_size, 0, null, null);
if(err < 0) {
print_error("clCreateBuffer", err);
}
err = clEnqueueReadBuffer(queue, res_buff, CL_TRUE, 0,
Sizeof.cl_float*64, Pointer.to(result), 0, null, null);
if(err < 0) {
print_error("clEnqueueReadBuffer", err);
}
for(int i = 0; i < 64; i++) {
System.out.printf("%f\n", result[i]);
}
clReleaseMemObject(res_buff);
clReleaseKernel(kernel);
clReleaseProgram(program);
clReleaseCommandQueue(queue);
clReleaseContext(context);
}
}出力.
0.000000 100.099998 200.000000 300.100006 10.010000 110.110001 210.009995 310.109985 20.000000 120.099998 220.000000 320.100006 30.010000 130.110001 230.009995 330.109985 1.001000 101.100998 201.001007 301.101013 11.011000 111.111000 211.011002 311.110992 21.000999 121.100998 221.001007 321.101013 31.011000 131.110992 231.011002 331.110992 2.000000 102.099998 202.000000 302.100006 12.010000 112.110001 212.009995 312.109985 22.000000 122.099998 222.000000 322.100006 32.009998 132.110001 232.009995 332.109985 3.001000 103.100998 203.001007 303.101013 13.011000 113.111000 213.011002 313.110992 23.000999 123.100998 223.001007 323.101013 33.011002 133.110992 233.011002 333.110992
出力を見て頂くと、グローバル空間が4x4x4、ローカル空間が2x2x2ということが確認できます。
詳細は「表：数学関数」（表B.103「表：数学関数」）を参照ください。
数学関数は組込み版のmath.hのようなものです。ベクトル型に対応しているので、math.hのアップグレード版と考えることもできます。math.hからの円滑な移行を可能するために設計された数学関数には、信号処理、統計、機械学習で頻繁に使うガンマ関数、三角関数、指数関数が提供されています。
関数定義.
gentype cos (gentype)
gentype sin (gentype)
gentype tan (gentype)
gentype sincos (
gentype x,
__global gentype *cosval)
サンプル.
float4 data = (float4)(M_PI_F,M_PI_2_F,M_PI_4_F,0.0f);
printf("data: %#v4f\n", data);
float4 res_cos = cos(data);
float4 res_sin = sin(data);
float4 res_sin_cos = res_sin/res_cos;
float4 res_tan = tan(data);
printf("cos(vecotr data) = %#v4f\n",res_cos);
printf("sin(vector data) = %#v4f\n",res_sin);
printf("sin/cos = %#v4f\n",res_sin_cos);
printf("tan = %#v4f\n", res_tan);
printf("cos(pi) = %f\n", cos(M_PI_F));
printf("cos(pi/2) = %f\n", cos(M_PI_2_F));
printf("cos(pi/4) = %f\n", cos(M_PI_4_F));
float angle = 45.0f;
float radians = angle * M_PI_F / 180.0f;
printf("45 degrees to radians, then cosine(r) = %f\n", cos(radians));
printf("pi/4 radians to degree %f\n", degrees(radians));
M_PI_F等の定数については「表：浮動小数点マクロ」（表：浮動小数点マクロ）と「表：浮動小数点との互換性」（表B.106「表：浮動小数点との互換性」）と「表：浮動小数点定数」（表B.107「表：浮動小数点定数」）を参照ください。
出力.
data: 3.141593,1.570796,0.785398,0.000000 cos(vecotr data) = -1.000000,-0.000000,0.707107,1.000000 sin(vector data) = -0.000000,1.000000,0.707107,0.000000 sin/cos = 0.000000,-22877334.000000,1.000000,0.000000 tan = 0.000000,-22877334.000000,1.000000,0.000000 cos(pi) = -1.000000 cos(pi/2) = -0.000000 cos(pi/4) = 0.707107 45 degrees to radians, then cosine(r) = 0.707107 pi/4 radians to degree 45.000000
gentypeは「float, float2, float3, float4, float8, float16, double, double2, double3, double4, double8 or double16」のいずれかを指す。
双曲線余弦（hyperbolic cosine）、双曲線正弦(hyperbolic sine)、双曲線正接（Hyperbolic arc tangent)を計算するには以下の関数を使います。
関数定義.
gentype cosh (gentype) gentype sinh (gentype) gentype tanh (gentype)
サンプル.
float res_sinh = sinh(2.0f);
float res_cosh = cosh(2.0f);
float res_tanh = tanh(2.0f);
printf("sinh=%f, cosh=%f, tanh=%f\n", res_sinh, res_cosh, res_tanh);
出力.
sinh=3.626860, cosh=3.762196, tanh=0.964028
逆双曲線余弦(Inverse hyperbolic cosine)、逆双曲線正弦(Inverse hyperbolic sine)、逆双曲線正接（Inverse Hyperbolic arc tangent)を計算するには以下の関数を使います。
関数定義.
gentype acosh (gentype) gentype asin (gentype) gentype atanh (gentype)
サンプル.
float res_asinh = asinh(3.626860f);
float res_acosh = acosh(3.762196f);
float res_atanh = atanh(0.964028f);
printf("asinh=%f, acosh=%f, atanh=%f\n", res_asinh, res_acosh, res_atanh);
出力.
sinh=3.626860, cosh=3.762196, tanh=0.964028
xのy乗を計算するには、pow関数や、pown関数を使います。pownのy引数は整数型です。
gentype pow(gentype x, gentype y) gentype pown(gentype x, intn y)
数値の丸めをする関数もあります。
gentype rint (gentype)
最近接偶数への丸め (round to the nearest even) モードで整数値への丸めを行います。
gentype ceil (gentype)
正の無限大への丸めモードで天井値を計算。
gentype round (gentype x)
丸め方向と関係なくxに最も近い整数値への丸め(中間値は０より遠い方向)を行います。
gentype trunc (gentype)
0への丸め (rounding toward zero)モードで整数への丸めを計算します。
サンプル.
float2 tmp = (float2)(1.234f, 2.654f);
float2 rint_out = rint(tmp);
printf("rint: %#v2f\n", rint_out);
float2 round_out = round(tmp);
printf("round: %#v2f\n", round_out);
float2 trunc_out = trunc(tmp);
printf("trunc: %#v2f\n", trunc_out);
float2 floor_out = floor(tmp);
printf("floor: %#v2f\n", floor_out);
float2 ceil_out = ceil(tmp);
printf("ceil: %#v2f\n", ceil_out);
出力.
rint: 1.000000,3.000000 round: 1.000000,3.000000 trunc: 1.000000,2.000000 floor: 1.000000,2.000000 ceil: 2.000000,3.000000
詳細は「表：half型、native型の数学関数」（<<half_maths>）を参照ください。
half型とnative型に適合した数学関数です。half_func_name、native_func_nameというように、 half_ か native_ が関数の先頭文字となります。
half_が接頭にある場合は半精度浮動小数点を使い計算をします。
gentype half_cos (gentype x)
この関数は余弦（cosine）の計算をし、引数xの範囲は -216 ... +216 とします。
半精度浮動小数点関数を使うのは、深層学習が流行したということもありますが、ハードウェアがニーズに対応した側面もあります。半精度浮動小数点は一般に単精度や倍精度の演算よりも高速ですが、ハードが対応しない場合はパフォーマンスが向上するとは限りません。
gentypeは「float, float2, float3, float4, float8, float16」のいずれかを指します。
nativeが窃盗にある関数はハードウェアの特殊演算ユニットに１対１または１対多のマッピングを行ないます。三角関数、指数関数、対数やべき乗の計算で高速な演算を行なうことができます。精度についてはOpenCLの実装によります。
gentype native_cos (gentype x)
余弦（cosine）の計算をしますが、引数xの範囲は実装システムに依存します。
詳しくは「表：浮動小数点境界値」（表B.105「表：浮動小数点境界値」）を参照ください。
詳しくは「表：浮動小数点マクロ」（表：浮動小数点マクロ）と「表：浮動小数点との互換性」（表B.106「表：浮動小数点との互換性」）と「表：浮動小数点定数」（表B.107「表：浮動小数点定数」）を参照ください。
詳しくは「表：倍精度型マクロ」（表：倍精度型マクロ）と「表：倍精度型の互換性」（表B.108「表：倍精度型の互換性」）と「表：倍精度定数」（表B.109「表：倍精度定数」）を参照ください。
Copyright 2018-2019, by Masaki Komatsu