ATI Stream SDK 2.0 beta と OpenCL

まだ未対応らしく Windows7 にはうまくインストールできませんでした。

ATI Stream Software Development Kit (SDK) v2.0 Beta Program

インストーラが正常終了しても Program Files の ATI Stream フォルダが
空のままです。
とりあえず Vista で展開したのちパスを通せば使えなくはないようです。
よく見ると OpenCL は GPU 未対応で CAL も無くなっています。
Intel CPU (+GeForce) でも動きました。

CL_PLATFORM_PROFILE    = FULL_PROFILE
CL_PLATFORM_VERSION    = OpenCL 1.0 ATI-Stream-v2.0-beta2
CL_PLATFORM_NAME       = ATI Stream
CL_PLATFORM_VENDOR     = Advanced Micro Devices
CL_PLATFORM_EXTENSIONS = 

CL_DEVICE_TYPE = 2
CL_DEVICE_VENDOR_ID = 4098
CL_DEVICE_MAX_COMPUTE_UNITS = 8
CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS = 3
CL_DEVICE_MAX_WORK_ITEM_SIZES = 1024 1024 1024
CL_DEVICE_MAX_WORK_GROUP_SIZE = 1024
CL_DEVICE_EXECUTION_CAPABILITIES = 1

Core i7 の結果です。
CL_DEVICE_TYPE の 2 は CPU (CL_DEVICE_TYPE_CPU) を表しています。
CL_DEVICE_MAX_COMPUTE_UNITS は 8個。
Core2 Duo で走らせると CL_DEVICE_MAX_COMPUTE_UNITS = 2 となります。

Device Type は下記の通り 3種類定義されています。

2 CL_DEVICE_TYPE_CPU
4 CL_DEVICE_TYPE_GPU
8 CL_DEVICE_TYPE_ACCELERATOR

Compute Shader は GPU 上で走るプログラムを共通化するのが目的でした。
OpenCL は CPU/GPU やその他を含めた、より幅広いプロセッサを対象として
作られているようです。

Direct3D の Compute Shader も CPU でエミュレーションできれば同じですが
現状だと WARP (10.1) も対応しておらず、Reference でしか動作しません。
エミュレーションできても GPU, CPU 混在しての利用は想定されていないといって
良いでしょう。

CL_DEVICE_EXECUTION_CAPABILITIES は CL_EXEC_KERNEL のみセットされています。
CL_EXEC_NATIVE_KERNEL が無いので、clEnqueueNativeKernel()
は実行出来ないようです。

バッファの転送や実行は Command Queue を生成して非同期に監視します。
GPU の描画操作に似ています。
異なるのは Out of order で実行可能なこと。
そのままでは実行順の保証が出来ませんが、その代わり各コマンドに依存関係を
設定できます。特定のコマンドの完了を待ってから実行できるわけです。

Direct3D11 でも DeviceContext が分離され、Command List (Queue) に
蓄えられる GPU 命令と、そうでないものを区別できるようになっています。
ただこちらは、複数のスレッドで作成したコマンドを単一のコンテキストで実行する
ための仕組みです。コマンドバッファ内の実行順は決まっているため、
スケジューリングは呼び出す側に委ねられています。

OpenCL の kernel の記述言語は glsl とは別もので、より C言語に近いものです。
long は 64bit。vector も 16個まで扱えるようです。
4コンポーネントを超えた場合、各要素は16進数指定で s0, s1,~ sf といった表記
になっています。

D3D の Compute Shader はデバイスやリソース管理等を D3D のコンポーネントに
委ねています。Direct3D の一部であって言語も hlsl そのまま同じものでした。

OpenCL の場合完全に独立しており、リソース管理も言語も OpenCL の世界で作られています。
コマンドをいくつか拡張してレンダリング向け補助機能を追加すれば、新しい
3D API セットができるのではないかと思うくらい。
その分 ATI CAL とか、独自 SDK と比べると複雑になった印象を受けるかもしれません。
リソースは描画と共有可能で、D3D や OpenGL からも OpenCL のバッファを作成
出来るようになっています。

const int THREAD_SIZE= 32;
cl_int	  status= 0;
cl_uint	  nums;
cl_platform_id	pid= 0;
cl_device_id	device= 0;

status= clGetPlatformIDs( 1, &pid, &nums );
status= clGetDeviceIDs( pid, CL_DEVICE_TYPE_CPU, 1, &device, &nums );

cl_context  context= clCreateContextFromType( NULL, CL_DEVICE_TYPE_CPU, NULL, NULL, &status );
cl_command_queue   command= clCreateCommandQueue( context, device, CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE, &status );

cl_mem	mem= clCreateBuffer( context, CL_MEM_READ_WRITE, THREAD_SIZE * sizeof(cl_float), NULL, &status );

const char*   src=
    "__kernel void main( __global float* op )\n"
    "{\n"
    "   unsigned int   tid= get_global_id(0);\n"
    "   op[tid]= tid * tid;\n"
    "}";
size_t	srcsize= strlen( src );
cl_program  prog= clCreateProgramWithSource( context, 1, &src, &srcsize, &status );

status= clBuildProgram( prog, 0, NULL, NULL, NULL, NULL );

cl_kernel  kernel= clCreateKernel( prog, "main", &status );

status= clSetKernelArg( kernel, 0, sizeof(cl_mem), &mem );

size_t	gthread= THREAD_SIZE;
size_t	lthread= 1;
cl_event   event;
status= clEnqueueNDRangeKernel( command, kernel, 1, NULL, >hread,