スカラ型の中でも、half型は特殊ケースのため補足します。IEEE 754-2008規格という半精度浮動小数点演算の規格に準拠したデータ型でして、ここで詳細は記述しません。

IEEE 754規格は浮動小数点を（２進数）ビットで表すために以下のような計算式を示しています。

表したい浮動小数点値 = 符号 * 2^指数 * 仮数

数式とするのであれば以下のように表すことができます。

各要素は以下の２進数のビット数で表します。

指数部は下駄履き表現とし、バイアスは15となります。

例えば、指数部分の計算をする場合、指数からバイアス（又はエキセス）を下記のように減じます。

仮数部の最上位ビットは、指数部が00000の場合は0とし、それ以外（無限値となる11111を除外）を1とします。

この計算のより詳細な記述を知りたい方は、IEEE754規格を参照頂くようお願い致します。

OpenCL-Cではhalf型はあまり使う機会はないかと思います。half型で計算をするハードウェアがないため、非効率な処理になる可能性が高いからです。

OpenCL 1.2以降では、half値のバッファオブジェクトを指すポインタを宣言する場合に使います。

floatは単精度浮動小数点（32bit）のスカラデータ型です。原則としてIEEE 754に準拠します。

本書を執筆している段階では、ハードウェアの演算ユニットはfloatであればサポートしているので、パフォーマンスを最適化する場合は、ハードの仕様がfloat以外をサポートしない場合は、floatをベースのデータ型としてアルゴリズムの設計を行うことを推奨します。

floatの各要素は以下の２進数のビット数で表します。

最大値は(1)、最小値は(2)で表すことができます。

64ビットの倍精度浮動小数点型、doubleはOpenCL規格ではオプション機能であり、多くのGPUデバイスでサポートしていません。

本書で検証機として利用しているMac Miniでdoubleがサポートされているか確認してみました。下記のように「FP_64」が定義されているか検証コードを書いてみました。

/* OpenCL 1.2では手動設定は不要 */
#ifdef FP_64
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#endif

__kernel void helloworld(float f)
{
#ifdef FP_64
        printf("Double precision FP supported by GPU.\n");
#else
        printf("Double precision FP not supported by GPU.\n");
#endif
}

実行可能ファイルの出力結果は以下のようになります。

$ ./helloworld
kernel completed.
Double precision FP not supported by GPU.

出力結果にあるようにMac Miniでは倍精度浮動小数点はサポートしていないことを確認できました。

double型の値は64ビット（２進数）で表します。

doubleについては、ハードウェアでサポートされていないこともありますが、より大きなメモリ領域を確保する必要が生じるため、使える機会は限定されます。

Table B.1, “表：スカラデータ型”では基本的なデータ型とは別にCL_DEVICE_ADDRESS_BITSという値に依存するデータ型が存在します。以下がそのデータ型です。

CL_DEVICE_ADDRESS_BITSはclGetDeviceInfo関数のparam_name引数にCL_DEVICE_ADDRESS_BITSを指定すると求めることができます。CL_DEVICE_ADDRESS_BITSの値は32か64のいずれかになります。

例えばsize_tのサイズはCL_DEVICE_ADDRESS_BITSが32の場合は、32bit、64の場合は64bitとなります。この作用は上記にあげた他３つのデータ型で共通しています。

IntelのGPUデバイスではメモリアドレス空間の最適化のために、最低でも32-bitのデータサイズでのアクセスが推奨されています。以下のデータ型は32ビットデータ型の例です。

intは32ビット、char4は8 x 4、ushort2は16 x 2ということで32ビットを構成します。

さらによりメモリを効率的に使うには以下のような４つの32-bitをまとめたベクトルデータ型が良いとされています。

これは4 x 4 = 16バイトに相当します。

最適化について追記すると、OpenCLアプリケーションでパフォーマンスをあげるには、ホストプログラム内でデータを初期化するのは定石といっていいでしょう。その際には、本項目で解説したベクトルデータ型はOpenCL-C専用なため、ホストAPI側の対応データ型を使う必要があります。

例えば浮動小数点型、整数型のスカラデータをカーネル引数として渡すのであれば以下のようにホストプログラム内で宣言します。

ただしスカラデータ型の場合は、floatやintのままでもコンパイルは通ります。OpenCLのホストAPIの組込みデータ型が役に立つのは、ベクトルとしてデータを初期化する場合です。

OpenCL-C言語のfloatn、intnに対応するのは以下の２つです。

このようにベクトル型にせよスカラ型にせよ、ホスト側とOpenCL-C（カーネル）の型の相互変換は直感的に行なうことができます。