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引数で指定したものが該当します。

import numpy as np
import pyopencl as cl

devices = [cl.get_platforms()[0].get_devices(cl.device_type.GPU)[0]]

ctx = cl.Context(devices)
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags

ary = np.arange(64).astype(np.uint32)
ary_mem = cl.Buffer(ctx, mf.USE_HOST_PTR, hostbuf=ary)

program = cl.Program(ctx, """
    __kernel void ndrange_id(__global write_only uint *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 * 100000 + global_id_1 * 10000 + global_id_2 * 1000;
       f += local_id_0 * 100 + local_id_1 * 10 + local_id_2 * 1;
       output[index] = f; //(15)
    }
    """).build()

kernel = cl.Kernel(program, name="ndrange_id")

output = np.zeros(64).astype(np.uint32)
out_mem = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY | cl.mem_flags.USE_HOST_PTR, hostbuf=output)
kernel.set_arg(0, out_mem)

global_size = (4, 4, 4)
local_size = (4, 4, 4)

cl.enqueue_nd_range_kernel(queue, kernel, global_work_size=global_size, local_work_size=local_size)

cl.enqueue_read_buffer(queue, mem=out_mem, hostbuf=output)

queue.finish()

print(output)

ホストプログラム側の設定は、64個のワークアイテムを3次元に分割するものとします。

出力. 

[     0 100100 200200 300300  10010 110110 210210 310310  20020 120120
 220220 320320  30030 130130 230230 330330   1001 101101 201201 301301
  11011 111111 211211 311311  21021 121121 221221 321321  31031 131131
 231231 331331   2002 102102 202202 302302  12012 112112 212212 312312
  22022 122122 222222 322322  32032 132132 232232 332332   3003 103103
 203203 303303  13013 113113 213213 313313  23023 123123 223223 323323
  33033 133133 233233 333333]

出力を見て頂くと、グローバル空間が4x4x4、ローカル空間が2x2x2ということが確認できます。

数学関数は組込み版の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));

出力. 

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

双曲線余弦（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_func_name、native_func_nameというように、*half_*か*native_*が関数の先頭文字となります。

half_が接頭にある場合は半精度浮動小数点を使い計算をします。

gentype half_cos (gentype x)

この関数は余弦（cosine）の計算をし、引数xの範囲は`-216 … +216`とします。

半精度浮動小数点関数を使うのは、深層学習が流行したということもありますが、ハードウェアがニーズに対応した側面もあります。半精度浮動小数点は一般に単精度や倍精度の演算よりも高速ですが、ハードが対応しない場合はパフォーマンスが向上するとは限りません。

nativeが窃盗にある関数はハードウェアの特殊演算ユニットに１対１または１対多のマッピングを行ないます。三角関数、指数関数、対数やべき乗の計算で高速な演算を行なうことができます。精度についてはOpenCLの実装によります。

gentype native_cos (gentype x)

余弦（cosine）の計算をしますが、引数xの範囲は実装システムに依存します。