OpenCLによるGPGPU入門②

OpenCLによるGPGPU入門①で説明したソースコードを貼っておく。

以下のコードでは、ファイルの更新状況によって、
オンラインコンパイルとオフラインコンパイルを自動で選択するようにしてある。

その時、ファイルの最終更新日時を取得するのに、Windows依存のAPIを使用している。

また置いてあるプロジェクトはVC++で作成したので、
他のOS(Windows以外のOS)では
改行コードが異なるため注意して欲しい。

以下コードと作成したプロジェクトのダウンロード。

個人的なメモ。
エラーログの表示で、
コンパイラから帰ってきたデータの表示で一部バグってたかもしれないのは、
windowsとlinux系では改行コードが異なるのが原因？
つまり、windowsでも正し動作するように改行コードを書きかえる必要があるかも。
「\n」→「\r\n」と書き換えるって事。
……ところでまだバグってたかどうかは……まあ、今度やる機会があったら確認しようorz。



コード
/* Visual Studio の設定 プロジェクト(P) →[プロジェクト名]のプロパティ(P)... Alt+F7 →構成プロパティ→C/C++→全般→追加のインクルード ディレクトリ に、「$(AMDAPPSDKROOT)\include」を追加 →構成プロパティ→リンカ―→全般→追加のライブラリ ディレクトリ に、「$(AMDAPPSDKROOT)\lib\x86」を追加 →構成プロパティ→リンカ―→入力→追加の依存ファイル に、「OpenCL.lib」を追加するか、 ソースコード中に#includeなどと一緒に 「#pragma comment(lib,"OpenCL.lib")」と記述しておく。 */ /* サンプルの場所は、 「C:\Users\＜ユーザー名＞\AMD APP SDK\2.9」 */ //#include <stdafx.h>//プリコンパイル済みヘッダ //たぶん「stdio.h」などのインクルードファイルを既にコンパイルして、コンパイル時間を短縮できる。今回は使用しない。 #include <stdio.h> #include <tchar.h> #include <windows.h> //system("pause");に必要 #include <CL/cl.h> #pragma comment(lib,"OpenCL.lib") //プロトタイプ宣言 cl_platform_id* GetPlatformIDs( cl_uint* num_platforms ); //プラットフォームのIDを取得する関数 cl_device_id* GetDeviceIDs(cl_platform_id* platform_ID , cl_device_info param , cl_uint* num_devices ); //デバイスのIDを取得する関数 void PrintErrorCode( const cl_int ErrorCode ); int PrintError( cl_int ErrorCode , char* PrintStr ); int PrintBuildLog( const cl_program program, const cl_device_id device ); //ファイルの最終更新日時を取得して返す関数 int Get_ftLastWriteTime_SYSTEMTIME( LPCSTR FilePath , SYSTEMTIME* SysTimeBuf ); int Get_ftLastWriteTime_FILETIME ( LPCSTR FilePath , FILETIME* FileTimeBuf ); int ReadALL( char* FilePath , char** ReadData , fpos_t* fsize );//ファイルのデータを全てメモリに読み込む関数 int GetBinariesANDSaveToFile( cl_program program , char *FilePath); int main( void ){ //プラットフォームのIDを取得する。(使い終わったら解放する。) cl_uint num_platforms=0; cl_platform_id* platform_IDs = GetPlatformIDs(&num_platforms); if( platform_IDs==NULL ){ return -1; } //デバイスのIDを取得する。(使い終わったら解放する。) cl_uint num_devices = 0; cl_device_id* device_IDs = GetDeviceIDs( &platform_IDs[0] , CL_DEVICE_TYPE_GPU , &num_devices ); if( device_IDs==NULL ){ return -1; } //コンテキストを作成。(使い終わったら解放する。) //clCreateContext()関数、或いは、clCreateContextFromType()関数、を用いてコンテキストを作成する。 cl_int ErrorCode; cl_context context = clCreateContext( NULL, //プロパティ //指定できる値は「cl.h」を「cl_context_properties」で検索すると、「CL_CONTEXT_PLATFORM」「CL_CONTEXT_INTEROP_USER_SYNC」 1, //デバイスの数 &device_IDs[0], //コンテキストを作成するデバイスのID。 NULL, //コンテキスト内でエラーが発生した場合に呼び出されるコールバック関数のポインタ NULL, //コールバック関数に引数として渡されるデータ &ErrorCode //関数の実行結果 ); if( PrintError( ErrorCode , "clCreateContext()" ) == -1 ){ return -1; } //■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■ //ファイルの最終更新日時の確認と、ロードするプログラムファイルの選択 //この部分には環境に依存した関数を使っている。(Win7 x64 Visual C++ 2012) /* //「Kernel.cl」 の最終更新日時を取得する。 LPCSTR FilePath_cl="./Kernel.cl"; SYSTEMTIME SysTimeBuf_cl; if( Get_ftLastWriteTime_SYSTEMTIME( FilePath_cl , &SysTimeBuf_cl ) == -1 ){return -1;} printf("%04d年%02d月%02d日 : %02d時%02d分%02d秒\n", SysTimeBuf_cl.wYear , SysTimeBuf_cl.wMonth , SysTimeBuf_cl.wDay , SysTimeBuf_cl.wHour , SysTimeBuf_cl.wMinute , SysTimeBuf_cl.wSecond ); //「Kernel.clbin」の最終更新日時を取得する。 LPCSTR FilePath_clbin="./Kernel.clbin"; SYSTEMTIME SysTimeBuf_clbin; if( Get_ftLastWriteTime_SYSTEMTIME( FilePath_clbin , &SysTimeBuf_clbin ) == -1 ){return -1;} printf("%04d年%02d月%02d日 : %02d時%02d分%02d秒\n", SysTimeBuf_clbin.wYear , SysTimeBuf_clbin.wMonth , SysTimeBuf_clbin.wDay , SysTimeBuf_clbin.wHour , SysTimeBuf_clbin.wMinute , SysTimeBuf_clbin.wSecond ); */ //二つの時刻を比較する。 //「バイナリファイル(コンパイル済みのファイル)の更新日時 >= ソースファイルの更新日時」: (コンパイルの必要が無いので)オフラインコンパイル //「バイナリファイル(コンパイル済みのファイル)の更新日時 < ソースファイルの更新日時」: (コンパイルの必要があるので)オンラインコンパイル FILETIME FileTimeBuf_cl; if( Get_ftLastWriteTime_FILETIME( "./Kernel.cl" , &FileTimeBuf_cl ) ){printf("Kernel.cl file is not exist\n");return -1;} FILETIME FileTimeBuf_clbin; int result_GFtime = Get_ftLastWriteTime_FILETIME( "./Kernel.clbin" , &FileTimeBuf_clbin );//ファイルが存在しなかった場合(result == -1の場合)は、必然的にオンラインコンパイル。 // 値 意味 // -1 第 1 のファイル時刻が第 2 のファイル時刻より小さくなっています。 //(コンパイルの必要が無いので)オフラインコンパイル // 0 第 1 のファイル時刻と第 2 のファイル時刻が等しくなっています。 //(コンパイルの必要があるので)オンラインコンパイル // +1 第 1 のファイル時刻が第 2 のファイル時刻より大きくなっています。 //(コンパイルの必要があるので)オンラインコンパイル cl_program program; if( CompareFileTime( &FileTimeBuf_cl , &FileTimeBuf_clbin ) == -1 && result_GFtime != -1 ){ //(コンパイルの必要が無いので)オフラインコンパイル //コンパイル済みバイナリプログラムからのプログラムの起動 //バイナリファイルの読み込み char *KernelBinary=NULL; fpos_t fsize=NULL; ReadALL( "./Kernel.clbin" , &KernelBinary , &fsize ); //プログラムオブジェクトの作成 cl_int ErrorCode_binary; program = clCreateProgramWithBinary( context, //プログラムオブジェクトを作成するコンテキスト 1, //ソースコードの文字列を格納する配列の要素数 &device_IDs[0], //プログラムを実行するデバイスのIDの配列 (const size_t*)&fsize, //ソースコードの文字列の文字数を格納する配列。各要素に0を指定するとその要素がNULLで終端されていると判断される。配列自体をNULLとすると、全ての、ソースコードの文字列を格納する配列が、NULLで終端されていると判断される。 (const unsigned char**)&KernelBinary,//バイナリプログラムへの二次元配列 &ErrorCode_binary, //バイナリデータについてのステータスへの配列 &ErrorCode //関数の実行結果 ); if( PrintError( ErrorCode , "clCreateProgramWithBinary()" ) == -1 ){ return -1; } if( PrintError( ErrorCode_binary , "clCreateProgramWithBinary()" ) == -1 ){ return -1; } free( KernelBinary );//メモリの解放 }else{ //(コンパイルの必要があるので)オンラインコンパイル //カーネルプログラムの読み込み //(「Kernel.cl」を文字列として読み込む。) char *KernelSource_str=NULL; fpos_t fsize=NULL; ReadALL( "./Kernel.cl" , &KernelSource_str , &fsize ); //プログラムオブジェクトの作成 //コンテキストを指定して、カーネルプログラムのソースコードの文字列から、プログラムオブジェクトを作成する。 //clCreateProgramWithBuiltInKernels() //類似の関数 program = clCreateProgramWithSource( context, //プログラムオブジェクトを作成するコンテキスト 1, //ソースコードの文字列を格納する配列の要素数 (const char**)&KernelSource_str,//ソースコードの文字列を格納する配列 (const size_t*)&fsize,//NULL, //ソースコードの文字列の文字数を格納する配列。各要素に0を指定するとその要素がNULLで終端されていると判断される。配列自体をNULLとすると、全ての、ソースコードの文字列を格納する配列が、NULLで終端されていると判断される。 &ErrorCode //関数の実行結果 ); if( PrintError( ErrorCode , "clCreateProgramWithSource()" ) == -1 ){ return -1; } free( KernelSource_str );//メモリの解放 } //■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■ //プログラムオブジェクトのビルド //一部のサイトや参考書で、「clCreateProgramWithBinary()関数」を使う場合、この関数(「clBuildProgram()」)は不要であるとあるが、実際問題としてそれでは動かないので、これは誤りだと考えられる。……ここで、保存したのは本当に実行可能形式のファイルだったのか?という疑問が起こる。あるいは、どちらの場合でもこの関数の実行は必要なのか? //デバイスを指定して、プログラムオブジェクトをビルドする。 ErrorCode = clBuildProgram( program, //プログラムオブジェクト 1, //プログラムを実行するデバイスのIDの配列、に格納されているIDの要素数 &device_IDs[0], //プログラムを実行するデバイスのIDの配列 NULL, //コンパイラに渡すオプション NULL, //ビルド終了時に呼び出されるコールバック関数 NULL //コールバック関数に引数として渡されるデータ ); if( ErrorCode != CL_SUCCESS ){ printf("clBuildProgram() was failed. return == "); PrintErrorCode( ErrorCode ); if( PrintBuildLog( program , device_IDs[0] ) == -1 ){ return -1; } system("pause"); return -1; } //■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■ if( CompareFileTime( &FileTimeBuf_cl , &FileTimeBuf_clbin ) == -1 && result_GFtime != -1 ){ //if文の中に「clBuildProgram」と一緒にまとめてもいいのだが、辞めておく。あと、可読性は下がるが、「CompareFileTime」関数の戻り値を記憶しておいて使ってもいい。 //(コンパイルの必要が無いので)オフラインコンパイル //この時は何もしない。 }else{ //(コンパイルの必要があるので)オンラインコンパイル //コンパイルが完了したバイナリデータを保存しておく。 if( GetBinariesANDSaveToFile( program , "./Kernel.clbin" ) == -1 ){return -1;} } //■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■■ //カーネルオブジェクトの作成 //ビルドされたプログラムオブジェクトから、カーネルを作成する。 //カーネルを作成する事によって、カーネルプログラムをC言語の関数のように使えるようになる……と思う。 cl_kernel my_kernel = clCreateKernel( program, //プログラムオブジェクト "addVector", //カーネル関数の名前 &ErrorCode //関数の実行結果 ); if( ErrorCode != CL_SUCCESS ){ printf("clCreateKernel() was failed. return == "); PrintErrorCode( ErrorCode ); system("pause"); return -1; } //in1、in2に計算させるための初期値をセット。outは初期化 float in1[1000]; float in2[1000]; float out[1000]; for(int i=0; i<1000; i++ ){ in1[i] = (float)i * 100.0f; in2[i] = (float)i / 100.0f; out[i] = 0.0f; } //メモリオブジェクトの作成。 //(基本的に、ホスト≒CPU、カーネル≒GPU) cl_mem memobj_in1 = clCreateBuffer( context, //メモリオブジェクト作成するコンテキスト CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, //フラグ sizeof(in1), //カーネル側に確保するメモリのサイズ in1, //ホスト上のメモリへのポインタ &ErrorCode //関数の実行結果 ); if( ErrorCode != CL_SUCCESS ){ printf("clCreateBuffer() was failed. return == "); PrintErrorCode( ErrorCode ); system("pause"); return -1; } //「cl.h」の「cl_mem_flags - bitfield」より、フラグに指定できる値。 //CL_MEM_READ_WRITE //(カーネルから)読み書き可能(デフォルト) //CL_MEM_WRITE_ONLY //(カーネルから)書き込み専用 //CL_MEM_READ_ONLY //(カーネルから)読み込み専用 //CL_MEM_USE_HOST_PTR //ホストとカーネルのメモリを同期する。(CPUとGPUのメモリを同期する。)(第4引数:host_ptr、が必要。) //CL_MEM_ALLOC_HOST_PTR //ホストからアクセス可能な場所にメモリを確保(たぶんGPU上のグローバルメモリー?……あるいはCPUのメインメモリー。) //CL_MEM_COPY_HOST_PTR //カーネル側にメモリを確保し、ホストのデータで初期化。(第4引数:host_ptr、が必要。) cl_mem memobj_in2 = clCreateBuffer( context, //メモリオブジェクト作成するコンテキスト CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, //フラグ sizeof(in2), //カーネル側に確保するメモリのサイズ in2, //ホスト上のメモリへのポインタ &ErrorCode //関数の実行結果 ); if( PrintError( ErrorCode , "clCreateBuffer()" ) == -1 ){ return -1; } cl_mem memobj_out = clCreateBuffer( context, //メモリオブジェクト作成するコンテキスト CL_MEM_WRITE_ONLY, //フラグ sizeof(out), //カーネル側に確保するメモリのサイズ NULL, //ホスト上のメモリへのポインタ &ErrorCode //関数の実行結果 ); if( PrintError( ErrorCode , "clCreateBuffer()" ) == -1 ){ return -1; } //カーネル関数の引数を設定。 ErrorCode = clSetKernelArg( my_kernel, //カーネルオブジェクト 0, //対応するカーネル関数の引数の位置を指定する。(第0引数～第X引数まで) sizeof(cl_mem), //セットする引数のサイズ。配列の場合は配列一つ分のサイズ。 (const void *)&memobj_in1 //セットする引数へのポインタ ); if( PrintError( ErrorCode , "clSetKernelArg()" ) == -1 ){ return -1; } ErrorCode = clSetKernelArg( my_kernel, //カーネルオブジェクト 1, //対応するカーネル関数の引数の位置を指定する。(第0引数～第X引数まで) sizeof(cl_mem), //セットする引数のサイズ。配列の場合は配列一つ分のサイズ。 (const void *)&memobj_in2 //セットする引数へのポインタ ); if( PrintError( ErrorCode , "clSetKernelArg()" ) == -1 ){ return -1; } ErrorCode = clSetKernelArg( my_kernel, //カーネルオブジェクト 2, //対応するカーネル関数の引数の位置を指定する。(第0引数～第X引数まで) sizeof(cl_mem), //セットする引数の型のサイズ。つまり、配列の場合は配列一つ分のサイズ。 (const void *)&memobj_out //セットする引数へのポインタ ); if( PrintError( ErrorCode , "clSetKernelArg()" ) == -1 ){ return -1; } //コマンドキューの作成。(使い終わったら解放する。) //(取得したデバイスに対してコマンドキューを作成する。)(コマンドキューはコンテキストに対して作成される。) //(コマンドキューからデバイスにコマンドを送れる。) cl_command_queue gpu_command_queue = clCreateCommandQueue( context, //コマンドキューを作成するコンテキスト device_IDs[0], //コマンドを送るデバイス NULL, //プロパティ //指定できる値は「cl.h」を「cl_command_queue_properties - bitfield」で検索すると、「CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE」「CL_QUEUE_PROFILING_ENABLE」 &ErrorCode //関数の実行結果 ); if( PrintError( ErrorCode , "clCreateCommandQueue()" ) == -1 ){ return -1; } //カーネルの実行(命令の発行。Enqueue(順番待ちの列)に命令を入れる。) //clEnqueueNDRangeKernel()関数を簡略化した関数として、clEnqueueTask()関数がある。これは単一ワークアイテムの場合に使用できる。 const size_t global_work_offset[]={30}; //とりあえず「global_work_offset」と「global_work_size」の大小関係に関わらず、「clEnqueueNDRangeKernel()関数」はエラーを返さない。 const size_t global_work_size[] ={300};//タスクの総数 //global_work_sizeはlocal_work_sizeで割り切れる必要がある。(この時、global_work_offsetの値は影響しない。) const size_t local_work_size[] ={0}; //1グループあたりのタスクの数 //第6引数に直接0を指定した場合は、「global_work_size」が「local_work_size」で割り切れずともエラーを吐かないが、この配列を通して指定すると、エラーを吐く。 //「local_work_size」の指定が問題となるのは、ワークグループを分割した場合だが、おそらくワークグループを分割する事によってローカルIDを使えるようになる事に、分割のメリットがある。ただ、インデックス空間の次元が増える事との違いが今一ピンと来ていないのでいいかげんな話だなぁ……。 //恐らくだけど、カーネル関数の部分で、「get_local_id()関数」とかを使うんじゃないかな?その時「get_work_dim()」を使って(有効なのは0～(get_work_dim()-1)[次元])次元を所得するとかしないとか? ErrorCode = clEnqueueNDRangeKernel( gpu_command_queue, //命令を発行するコマンドキューを指定 my_kernel, //ビルドされたプログラムオブジェクトがセットされた、カーネルオブジェクトを指定 1, //カーネルを実行するインデックス空間の次元 global_work_offset,//NULL, //ワークアイテムのオフセットを配列として指定する。　　　 (この値から、OpenCL C言語部分において、「get_global_id()関数」よりidが発行される。) global_work_size, //　　　ワークアイテムの総数を配列として指定する。(並列化の数?)(この値まで、OpenCL C言語部分において、「get_global_id()関数」よりidが発行される。) //インデックス空間の各次元におけるワークアイテムの数を配列として指定する。 0,// local_work_size,//0, //ワークグループを分割する場合に、一つのワークグループに含まれるワークアイテムの数を指定する。??? 0, //event_wait_list(すぐ下の引数)に指定されるイベントオブジェクトの数 NULL, //ここに指定したイベントオブジェクトが終了するまで、コマンドの実行を待機する。 NULL //カーネルを実行した際に、ここに指定したイベントオブジェクトに対してイベントを発生させる(通知する?)。 ); if( PrintError( ErrorCode , "clEnqueueNDRangeKernel()" ) == -1 ){ return -1; } //メモリオブジェクトから結果を取得する。(データのリードコマンドをキューにセットする。) ErrorCode = clEnqueueReadBuffer( gpu_command_queue, //命令を発行するコマンドキューを指定 memobj_out, //データを読み込むメモリオブジェクト CL_BLOCKING, //このコマンドが完了するまで、この関数がブロックされるかどうかのフラグ。 0, //読み込み先ののデータの、先頭からのオフセットバイト数 sizeof(cl_float)*1000, //読み込むデータのサイズ out, //カーネルから読み込んだデータを格納する、ホスト側のバッファへのポインタ 0, //event_wait_list(すぐ下の引数)に指定されるイベントオブジェクトの数 NULL, //ここに指定したイベントオブジェクトが終了するまで、コマンドの実行を待機する。 NULL //カーネルを実行した際に、ここに指定したイベントオブジェクトに対してイベントを発生させる(通知する?)。 ); if( PrintError( ErrorCode , "clEnqueueReadBuffer()" ) == -1 ){ return -1; } //類似の関数 //clEnqueueReadBuffer()関数 : カーネル から ホスト へデータを転送(コピー)する。 //clEnqueueWriteBuffer()関数 : ホスト から カーネル へデータを転送(コピー)する。 //clEnqueueCopyBuffer()関数 : カーネル から カーネル へデータを転送(コピー)する。(メモリオブジェクト間のデータ転送(コピー)) //結果の表示(始めと終わりの一部のみ) printf(" in1 , in2 , out ( 検算結果 )\n"); for(int i=0; i<50; i++){ printf(" %12f , %12f , %12f ( %12f )\n",in1[i],in2[i],out[i],in1[i]+in2[i]); } printf("\n\n"); for(int i=950; i<1000; i++){ printf(" %12f , %12f , %12f ( %12f )\n",in1[i],in2[i],out[i],in1[i]+in2[i]); } system("pause"); //リソースの解放 clReleaseCommandQueue( gpu_command_queue ); //取得に失敗した場合は解放しなくていい……、のか？ clReleaseMemObject( memobj_out ); clReleaseMemObject( memobj_in2 ); clReleaseMemObject( memobj_in1 ); clReleaseKernel( my_kernel ); clReleaseProgram( program ); // clReleaseCommandQueue( gpu_command_queue ); //取得に失敗した場合は解放しなくていい……、のか？ clReleaseContext( context ); //取得に失敗した場合は解放しなくていい……、のか？ free( device_IDs ); free( platform_IDs ); return 0; } //プラットフォームのIDを取得する関数 cl_platform_id* GetPlatformIDs( cl_uint* num_platforms ){ //引数 // num_platforms : プラットフォームの数を返す。 // //戻り値 // NULL :失敗 // platform_IDs :成功(プラットフォームのIDを返す。) // //成功した場合は、終了時にメモリの解放が必要 cl_int result; //プラットフォームの総数を取得する。 result = clGetPlatformIDs(0, NULL, num_platforms); //成功すると利用可能なプラットフォームの数が返される。 if(result != CL_SUCCESS){ printf("When searching number of platforms\n"); //printf("clGetPlatformIDs() was failed. return == %d\n", result); printf("clGetPlatformIDs() was failed. return == "); PrintErrorCode( result ); system("pause"); return NULL; } //必要な量のメモリを確保。 cl_platform_id* platform_IDs = (cl_platform_id*)malloc( sizeof(cl_platform_id) * (*num_platforms) ); //プラットフォームIDの取得。 result = clGetPlatformIDs(*num_platforms, platform_IDs, NULL); if(result != CL_SUCCESS){ printf("When trying to get platform IDs\n"); //printf("clGetPlatformIDs() failed. return == %d\n", result); printf("clGetPlatformIDs() failed. return == "); PrintErrorCode( result ); free( platform_IDs ); system("pause"); return NULL; } return platform_IDs; } //デバイスのIDを取得する関数 cl_device_id* GetDeviceIDs(cl_platform_id* platform_ID , cl_device_info param , cl_uint* num_devices ){ //引数 // platform_ID : clGetPlatformIDs又は、GetPlatformIDsで取得したプラットフォームのIDを渡す。 // //戻り値 // NULL :失敗 // device_IDs :成功(デバイスのIDを返す。) // //成功した場合は、終了時にメモリの解放が必要 //引数のメモ(この引数で指定したデバイスのIDの数やIDを取得する。) // CL_DEVICE_TYPE_CPU // CL_DEVICE_TYPE_GPU // CL_DEVICE_TYPE_ACCELERATOR // CL_DEVICE_TYPE_DEFAULT // CL_DEVICE_TYPE_ALL cl_int result; //デバイスの数を取得する。 result = clGetDeviceIDs( *platform_ID , param , 0 , NULL , num_devices ); if( result != CL_SUCCESS ){ printf("When trying to get number of devices\n"); //printf("clGetDeviceIDs() was failed. return == %d\n", result); printf("clGetDeviceIDs() was failed. return == "); PrintErrorCode( result ); system("pause"); return NULL; } //必要な量のメモリを確保。 cl_device_id* device_IDs = (cl_device_id*)malloc( sizeof(cl_device_id) * (*num_devices) ); //デバイスのIDを取得する。 result = clGetDeviceIDs( *platform_ID , param , *num_devices , device_IDs , 0 ); if( result != CL_SUCCESS ){ printf("When trying to get device IDs\n"); //printf("clGetDeviceIDs() failed. return == %d\n", result); printf("clGetDeviceIDs() was failed. return == "); PrintErrorCode( result ); free( device_IDs ); system("pause"); return NULL; } return device_IDs; } //エラーコードの表示 void PrintErrorCode( const cl_int ErrorCode ){ //「cl.h」の「Error Codes」と照らし合わせて出力する。 switch( ErrorCode ){ case CL_SUCCESS: printf("CL_SUCCESS"); break; case CL_DEVICE_NOT_FOUND: printf("CL_DEVICE_NOT_FOUND"); break; case CL_DEVICE_NOT_AVAILABLE: printf("CL_DEVICE_NOT_AVAILABLE"); break; case CL_COMPILER_NOT_AVAILABLE: printf("CL_COMPILER_NOT_AVAILABLE"); break; case CL_MEM_OBJECT_ALLOCATION_FAILURE: printf("CL_MEM_OBJECT_ALLOCATION_FAILURE"); break; case CL_OUT_OF_RESOURCES: printf("CL_OUT_OF_RESOURCES"); break; case CL_OUT_OF_HOST_MEMORY: printf("CL_OUT_OF_HOST_MEMORY"); break; case CL_PROFILING_INFO_NOT_AVAILABLE: printf("CL_PROFILING_INFO_NOT_AVAILABLE"); break; case CL_MEM_COPY_OVERLAP: printf("CL_MEM_COPY_OVERLAP"); break; case CL_IMAGE_FORMAT_MISMATCH: printf("CL_IMAGE_FORMAT_MISMATCH"); break; case CL_IMAGE_FORMAT_NOT_SUPPORTED: printf("CL_IMAGE_FORMAT_NOT_SUPPORTED"); break; case CL_BUILD_PROGRAM_FAILURE: printf("CL_BUILD_PROGRAM_FAILURE"); break; case CL_MAP_FAILURE: printf("CL_MAP_FAILURE"); break; case CL_MISALIGNED_SUB_BUFFER_OFFSET: printf("CL_MISALIGNED_SUB_BUFFER_OFFSET"); break; case CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST: printf("CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST"); break; case CL_COMPILE_PROGRAM_FAILURE: printf("CL_COMPILE_PROGRAM_FAILURE"); break; case CL_LINKER_NOT_AVAILABLE: printf("CL_LINKER_NOT_AVAILABLE"); break; case CL_LINK_PROGRAM_FAILURE: printf("CL_LINK_PROGRAM_FAILURE"); break; case CL_DEVICE_PARTITION_FAILED: printf("CL_DEVICE_PARTITION_FAILED"); break; case CL_KERNEL_ARG_INFO_NOT_AVAILABLE: printf("CL_KERNEL_ARG_INFO_NOT_AVAILABLE"); break; //============================================= case CL_INVALID_VALUE: printf("CL_INVALID_VALUE"); break; case CL_INVALID_DEVICE_TYPE: printf("CL_INVALID_DEVICE_TYPE"); break; case CL_INVALID_PLATFORM: printf("CL_INVALID_PLATFORM"); break; case CL_INVALID_DEVICE: printf("CL_INVALID_DEVICE"); break; case CL_INVALID_CONTEXT: printf("CL_INVALID_CONTEXT"); break; case CL_INVALID_QUEUE_PROPERTIES: printf("CL_INVALID_QUEUE_PROPERTIES"); break; case CL_INVALID_COMMAND_QUEUE: printf("CL_INVALID_COMMAND_QUEUE"); break; case CL_INVALID_HOST_PTR: printf("CL_INVALID_HOST_PTR"); break; case CL_INVALID_MEM_OBJECT: printf("CL_INVALID_MEM_OBJECT"); break; case CL_INVALID_IMAGE_FORMAT_DESCRIPTOR: printf("CL_INVALID_IMAGE_FORMAT_DESCRIPTOR"); break; case CL_INVALID_IMAGE_SIZE: printf("CL_INVALID_IMAGE_SIZE"); break; case CL_INVALID_SAMPLER: printf("CL_INVALID_SAMPLER"); break; case CL_INVALID_BINARY: printf("CL_INVALID_BINARY"); break; case CL_INVALID_BUILD_OPTIONS: printf("CL_INVALID_BUILD_OPTIONS"); break; case CL_INVALID_PROGRAM: printf("CL_INVALID_PROGRAM"); break; case CL_INVALID_PROGRAM_EXECUTABLE: printf("CL_INVALID_PROGRAM_EXECUTABLE"); break; case CL_INVALID_KERNEL_NAME: printf("CL_INVALID_KERNEL_NAME"); break; case CL_INVALID_KERNEL_DEFINITION: printf("CL_INVALID_KERNEL_DEFINITION"); break; case CL_INVALID_KERNEL: printf("CL_INVALID_KERNEL"); break; case CL_INVALID_ARG_INDEX: printf("CL_INVALID_ARG_INDEX"); break; case CL_INVALID_ARG_VALUE: printf("CL_INVALID_ARG_VALUE"); break; case CL_INVALID_ARG_SIZE: printf("CL_INVALID_ARG_SIZE"); break; case CL_INVALID_KERNEL_ARGS: printf("CL_INVALID_KERNEL_ARGS"); break; case CL_INVALID_WORK_DIMENSION: printf("CL_INVALID_WORK_DIMENSION"); break; case CL_INVALID_WORK_GROUP_SIZE: printf("CL_INVALID_WORK_GROUP_SIZE"); break; case CL_INVALID_WORK_ITEM_SIZE: printf("CL_INVALID_WORK_ITEM_SIZE"); break; case CL_INVALID_GLOBAL_OFFSET: printf("CL_INVALID_GLOBAL_OFFSET"); break; case CL_INVALID_EVENT_WAIT_LIST: printf("CL_INVALID_EVENT_WAIT_LIST"); break; case CL_INVALID_EVENT: printf("CL_INVALID_EVENT"); break; case CL_INVALID_OPERATION: printf("CL_INVALID_OPERATION"); break; case CL_INVALID_GL_OBJECT: printf("CL_INVALID_GL_OBJECT"); break; case CL_INVALID_BUFFER_SIZE: printf("CL_INVALID_BUFFER_SIZE"); break; case CL_INVALID_MIP_LEVEL: printf("CL_INVALID_MIP_LEVEL"); break; case CL_INVALID_GLOBAL_WORK_SIZE: printf("CL_INVALID_GLOBAL_WORK_SIZE"); break; case CL_INVALID_PROPERTY: printf("CL_INVALID_PROPERTY"); break; case CL_INVALID_IMAGE_DESCRIPTOR: printf("CL_INVALID_IMAGE_DESCRIPTOR"); break; case CL_INVALID_COMPILER_OPTIONS: printf("CL_INVALID_COMPILER_OPTIONS"); break; case CL_INVALID_LINKER_OPTIONS: printf("CL_INVALID_LINKER_OPTIONS"); break; case CL_INVALID_DEVICE_PARTITION_COUNT: printf("CL_INVALID_DEVICE_PARTITION_COUNT"); break; //============================================= default: printf("Not Defined Error Code : %d",ErrorCode); break; } printf("\n"); } //エラーコードの表示関数 int PrintError( cl_int ErrorCode , char* PrintStr ){ if( ErrorCode != CL_SUCCESS ){ printf("%s was failed. return == ", PrintStr ); PrintErrorCode( ErrorCode ); system("pause"); return -1; } return 0; } //ビルドエラーのログを表示 int PrintBuildLog( const cl_program program, const cl_device_id device ){ cl_int ErrorCode; size_t required_buf_size=0; //必要とされるバッファのサイズを取得 ErrorCode = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_LOG, //取得する情報を指定 NULL, NULL, &required_buf_size ); if( ErrorCode != CL_SUCCESS ){ printf("clGetProgramBuildInfo() was failed. return == "); PrintErrorCode( ErrorCode ); system("pause"); return -1; } //必要な量のメモリを確保。 char* buf = (char*)malloc( sizeof(char) * (required_buf_size) ); //情報の取得 ErrorCode = clGetProgramBuildInfo( program, device, CL_PROGRAM_BUILD_LOG, //取得する情報を指定 required_buf_size, buf, &required_buf_size ); if( ErrorCode != CL_SUCCESS ){ printf("clGetProgramBuildInfo() was failed. return == "); PrintErrorCode( ErrorCode ); system("pause"); return -1; } printf("%s\n",buf); return 0; } //ファイルの最終更新日時を取得して返す関数 int Get_ftLastWriteTime_SYSTEMTIME( LPCSTR FilePath , SYSTEMTIME* SysTimeBuf ){ //FilePath : 情報を取得するファイルへのパスを指定 //SysTimeBuf : 取得したシステムタイムを返す。 //「FindFirstFile()関数」か「GetFileTime()関数」を使う。どちらもWin32 API //「FindFirstFile()関数」は使わない情報も多くて煩雑なだけなので、今回は見送ろうかと思ったが、余分にインクルードが必要そうなので、こっちを使う。 WIN32_FIND_DATA findData; HANDLE hFile = FindFirstFile(FilePath, &findData); if (hFile == INVALID_HANDLE_VALUE){ printf("FindFirstFile() was failed."); system("pause"); return -1; } FindClose(hFile); // hFindのクローズ FileTimeToLocalFileTime(&findData.ftLastWriteTime,&findData.ftLastWriteTime); //世界標準時(UTC) に基づく時間をローカルファイル時刻形式に変換する //メモリの無駄なので、同じ構造体に戻してやる。 FileTimeToSystemTime(&findData.ftLastWriteTime,SysTimeBuf);//FILETIME を SYSTEMTIME に変換する return 0; } //ファイルの最終更新日時を取得して返す関数 int Get_ftLastWriteTime_FILETIME( LPCSTR FilePath , FILETIME* FileTimeBuf ){ //FilePath : 情報を取得するファイルへのパスを指定 //SysTimeBuf : 取得したシステムタイムを返す。 //「FindFirstFile()関数」か「GetFileTime()関数」を使う。どちらもWin32 API //「FindFirstFile()関数」は使わない情報も多くて煩雑なだけなので、今回は見送ろうかと思ったが、余分にインクルードが必要そうなので、こっちを使う。 WIN32_FIND_DATA findData; HANDLE hFile = FindFirstFile(FilePath, &findData); if (hFile == INVALID_HANDLE_VALUE){ // printf("FindFirstFile() was failed."); // // system("pause"); // return -1; } FindClose(hFile); // hFindのクローズ FileTimeToLocalFileTime(&findData.ftLastWriteTime,FileTimeBuf); //世界標準時(UTC) に基づく時間をローカルファイル時刻形式に変換する //メモリの無駄なので、同じ構造体に戻してやる。 // FileTimeToSystemTime(&findData.ftLastWriteTime,SysTimeBuf);//FILETIME を SYSTEMTIME に変換する return 0; }

以下の関数 では、FilePathに指定したファイルをmallocに確保したメモリ上に全てコピーします。
関数内でmallocによりメモリーを確保するため、使用後はfree()によるメモリの解放が必要です。
ReadDataに読み込んだデータの配列を返します。
また、fsizeに読み込んっだデータサイズを返します。
//ファイルのデータを全てメモリに読み込む関数 int ReadALL( char* FilePath , char** ReadData , fpos_t* fsize ){ //FilePath : 読み込み先のファイルのパス //ReadData : 読み込んだデータへのポインタ //関数が成功した場合、メモリの解放が必要。 FILE *fp; fopen_s(&fp,FilePath,"rb"); //テキストモードで開く if( fp == NULL ){ printf("fopen() was failed"); system("pause"); return -1; } //ファイルサイズの取得 fseek(fp,0,SEEK_END); //ファイルポインタをファイルの最後に移動させる。//バイナリモードでは、fseek()関数で「SEEK_END」するのは、バイナリデータが終端文字を含む可能性があるため、動作不定とされているが、実際に開いているのはただのテキストであるため、終端文字は存在しえない。そのため、問題ない。 (*fsize)=ftell(fp); //ファイルサイズを取得する。(どちらでもOK) // fgetpos(fp_cl,fsize); //ファイルサイズを取得する。(どちらでもOK)// fgetpos()関数を使えば、2GBを超えるファイルでも扱える……らしい。 fseek(fp,0L,SEEK_SET); //ファイルポインタを先頭に戻す。 //ファイルサイズ分の領域を確保 *ReadData = (char*)malloc( (size_t)(*fsize)*sizeof(char) );//「clCreateProgramWithSource」の第4引数でサイズを指定するので、終端文字(NULL文字)の分のバァッファは確保しない。 //ファイルの読み込み (*fsize) = fread( *ReadData , sizeof(char) , (size_t)(*fsize) , fp );//このようにすればテキストモードで読み取っても上手く行く。(「\r\n」→「\n」と変換される分だけ読み取った文字の数が変化するので、修正が必要。) fclose( fp ); return 0; }



「オフラインコンパイル」する場合に、必要なデータを保存しておく。

たしか……
programを渡しておけば、
FilePathに指定したファイルに勝手に保存してくれる。
次回以降、コンパイルの必要が無ければ、FilePathに保存したファイルを読み込めば良い。
FilePathに指定したファイルが存在しない場合は、勝手に作るんだったか、自分で作る必要があるのかは……忘れた。
//ビルドしたコンパイル済みバイナリプログラムを取得し保存する関数 int GetBinariesANDSaveToFile( cl_program program , char *FilePath){ cl_int ErrorCode; /* //ビルドしたコンパイル済みバイナリプログラムを取得する。 //プログラムオブジェクトに関連づけられている、デバイスの数を取得する。(今回はどうせ一つのデバイスしか扱わないので不要ではある。) size_t num_program_devices; ErrorCode = clGetProgramInfo( program, CL_PROGRAM_NUM_DEVICES, sizeof(num_program_devices), &num_program_devices, NULL ); if( PrintError( ErrorCode , "clGetProgramInfo()" ) == -1 ){ return -1; } //必要とされるバッファのサイズを取得する。 size_t *binary_sizes = (size_t *)malloc( sizeof(size_t)*num_program_devices ); ErrorCode = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t)*num_program_devices, binary_sizes, NULL ); if( PrintError( ErrorCode , "clGetProgramInfo()" ) == -1 ){ return -1; } //必要な量のメモリを確保。 char ** binaries = (char **)malloc( sizeof(char *)*num_program_devices ); for(int i=0; i<(int)num_program_devices ;i++){ if(binary_sizes[i] != 0){ binaries[i] = (char *)malloc( sizeof(char)*binary_sizes[i] ); }else{ binaries[i] = NULL; } } //メモリの解放 free( binary_sizes ); //バイナリの読み込み ErrorCode = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(char *)*num_program_devices , binaries, NULL );//ここの第4引数は、ダブルポインタでアクセスされるので注意。//第3引数は、第4引数がポインタなので、当然ポインタのサイズを渡す。 if( PrintError( ErrorCode , "clGetProgramInfo()" ) == -1 ){ return -1; } //(ファイルへ書き込み) //メモリの解放 for(int i=0; i<(int)num_program_devices ;i++){ free( binaries[i] ); } free( binaries ); */ //ビルドしたコンパイル済みバイナリプログラムを取得する。(簡略版(デバイスの数が1で固定の場合を考える。)) //必要とされるバッファのサイズを取得する。 size_t binary_sizes; ErrorCode = clGetProgramInfo( program, CL_PROGRAM_BINARY_SIZES, sizeof(size_t)*1, &binary_sizes, NULL ); if( PrintError( ErrorCode , "clGetProgramInfo()" ) == -1 ){ return -1; } //必要な量のメモリを確保。 char* binary = (char*)malloc( sizeof(char) * (binary_sizes) ); //バイナリの読み込み ErrorCode = clGetProgramInfo( program, CL_PROGRAM_BINARIES, sizeof(char *)*1 , &binary, NULL );//ここの第4引数は、ダブルポインタでアクセスされるので注意。//第3引数は、第4引数がポインタなので、当然ポインタのサイズを渡す。 if( PrintError( ErrorCode , "clGetProgramInfo()" ) == -1 ){ return -1; } //ファイルへ書き込み FILE *fp_clbin; fopen_s(&fp_clbin,FilePath,"wb"); //バイナリモードで開く//テキストモードで開いた場合、もし「\n」などのデータがあった場合、「\r\n」に変換されて保存されるのでマズイ。 if( fp_clbin == NULL ){ printf("fopen() was failed"); system("pause"); return -1; } fwrite( binary , sizeof(char) , (size_t)binary_sizes , fp_clbin );//このようにすればテキストモードで読み取っても上手く行く。(「\r\n」→「\n」と変換される分だけ読み取った文字の数が変化するので、修正が必要。) fclose( fp_clbin ); //メモリの解放 free( binary ); return 0; }




Kernel.clファイル(実際にGPU上で動作するプログラム) /* Visual Studio の設定 そのままの状態では、単なるテキストエディタとして開かれるため、 行番号の表示と、コードの色分けをさせるために以下の設定を行う。 [ツール(T)]→[オプション(O)...] →[テキストエディタ] →[ファイル拡張子] で、 「 拡張子(T):cl 」 「エディター(I):Microsoft Visual C++」 として [追加(D)]→[OK]//……この時下の一覧から拡張子cl選択し、「適応」を一応押して、対象のファイルをエディタで再起動した。たぶん再起動だけでOK? 古いサイトで、[usertype.dat]を追加しないと、OpenCL特有の情報を色分けしてくれない、とあったが、 現在([AMD APP SDK]の[2.9])では、[usertype.dat]そのものが検索を掛けても見つからない上に、 ここまでの設定で、 [get_global_id()関数]など、OpenCL特有と思われる関数も正しく色分けされているので、 不要なものだと考える。 ↓↓ ……と思っていたのだが……。 */ __kernel void addVector( __global float *in1, __global float *in2, __global float *out ){ size_t index_x = get_global_id(0); //配列が、N==1,2,3次元の時に使う。 //グローバルID、時にはローカルIDも……使える？、によって計算をコントロールする。 //size_t index_y = get_global_id(1); //配列が、N== 2,3次元の時に使う。 //グローバルID、時にはローカルIDも……使える？、によって計算をコントロールする。 //size_t index_z = get_global_id(2); //配列が、N== 3次元の時に使う。 //グローバルID、時にはローカルIDも……使える？、によって計算をコントロールする。 out[index_x] = in1[index_x] + in2[index_x]; }


//実行結果
in1 , in2 , out ( 検算結果 ) 0.000000 , 0.000000 , 0.000000 ( 0.000000 ) 100.000000 , 0.010000 , 0.000000 ( 100.010002 ) 200.000000 , 0.020000 , 0.000000 ( 200.020004 ) 300.000000 , 0.030000 , 0.000000 ( 300.029999 ) 400.000000 , 0.040000 , 0.000000 ( 400.040009 ) 500.000000 , 0.050000 , 0.000000 ( 500.049988 ) 600.000000 , 0.060000 , 0.000000 ( 600.059998 ) 700.000000 , 0.070000 , 0.000000 ( 700.070007 ) 800.000000 , 0.080000 , 0.000000 ( 800.080017 ) 900.000000 , 0.090000 , 0.000000 ( 900.090027 ) 1000.000000 , 0.100000 , 0.000000 ( 1000.099976 ) 1100.000000 , 0.110000 , 0.000000 ( 1100.109985 ) 1200.000000 , 0.120000 , 0.000000 ( 1200.119995 ) 1300.000000 , 0.130000 , 0.000000 ( 1300.130005 ) 1400.000000 , 0.140000 , 0.000000 ( 1400.140015 ) 1500.000000 , 0.150000 , 0.000000 ( 1500.150024 ) 1600.000000 , 0.160000 , 0.000000 ( 1600.160034 ) 1700.000000 , 0.170000 , 0.000000 ( 1700.170044 ) 1800.000000 , 0.180000 , 0.000000 ( 1800.180054 ) 1900.000000 , 0.190000 , 0.000000 ( 1900.189941 ) 2000.000000 , 0.200000 , 0.000000 ( 2000.199951 ) 2100.000000 , 0.210000 , 0.000000 ( 2100.209961 ) 2200.000000 , 0.220000 , 0.000000 ( 2200.219971 ) 2300.000000 , 0.230000 , 0.000000 ( 2300.229980 ) 2400.000000 , 0.240000 , 0.000000 ( 2400.239990 ) 2500.000000 , 0.250000 , 0.000000 ( 2500.250000 ) 2600.000000 , 0.260000 , 0.000000 ( 2600.260010 ) 2700.000000 , 0.270000 , 0.000000 ( 2700.270020 ) 2800.000000 , 0.280000 , 0.000000 ( 2800.280029 ) 2900.000000 , 0.290000 , 0.000000 ( 2900.290039 ) 3000.000000 , 0.300000 , 3000.300049 ( 3000.300049 ) 3100.000000 , 0.310000 , 3100.310059 ( 3100.310059 ) 3200.000000 , 0.320000 , 3200.320068 ( 3200.320068 ) 3300.000000 , 0.330000 , 3300.330078 ( 3300.330078 ) 3400.000000 , 0.340000 , 3400.340088 ( 3400.340088 ) 3500.000000 , 0.350000 , 3500.350098 ( 3500.350098 ) 3600.000000 , 0.360000 , 3600.360107 ( 3600.360107 ) 3700.000000 , 0.370000 , 3700.370117 ( 3700.370117 ) 3800.000000 , 0.380000 , 3800.379883 ( 3800.379883 ) 3900.000000 , 0.390000 , 3900.389893 ( 3900.389893 ) 4000.000000 , 0.400000 , 4000.399902 ( 4000.399902 ) 4100.000000 , 0.410000 , 4100.410156 ( 4100.410156 ) 4200.000000 , 0.420000 , 4200.419922 ( 4200.419922 ) 4300.000000 , 0.430000 , 4300.430176 ( 4300.430176 ) 4400.000000 , 0.440000 , 4400.439941 ( 4400.439941 ) 4500.000000 , 0.450000 , 4500.450195 ( 4500.450195 ) 4600.000000 , 0.460000 , 4600.459961 ( 4600.459961 ) 4700.000000 , 0.470000 , 4700.470215 ( 4700.470215 ) 4800.000000 , 0.480000 , 4800.479980 ( 4800.479980 ) 4900.000000 , 0.490000 , 4900.490234 ( 4900.490234 ) 95000.000000 , 9.500000 , 0.000000 ( 95009.500000 ) 95100.000000 , 9.510000 , 0.000000 ( 95109.507813 ) 95200.000000 , 9.520000 , 0.000000 ( 95209.523438 ) 95300.000000 , 9.530000 , 0.000000 ( 95309.531250 ) 95400.000000 , 9.540000 , 0.000000 ( 95409.539063 ) 95500.000000 , 9.550000 , 0.000000 ( 95509.546875 ) 95600.000000 , 9.560000 , 0.000000 ( 95609.562500 ) 95700.000000 , 9.570000 , 0.000000 ( 95709.570313 ) 95800.000000 , 9.580000 , 0.000000 ( 95809.578125 ) 95900.000000 , 9.590000 , 0.000000 ( 95909.593750 ) 96000.000000 , 9.600000 , 0.000000 ( 96009.601563 ) 96100.000000 , 9.610000 , 0.000000 ( 96109.609375 ) 96200.000000 , 9.620000 , 0.000000 ( 96209.617188 ) 96300.000000 , 9.630000 , 0.000000 ( 96309.632813 ) 96400.000000 , 9.640000 , 0.000000 ( 96409.640625 ) 96500.000000 , 9.650000 , 0.000000 ( 96509.648438 ) 96600.000000 , 9.660000 , 0.000000 ( 96609.656250 ) 96700.000000 , 9.670000 , 0.000000 ( 96709.671875 ) 96800.000000 , 9.680000 , 0.000000 ( 96809.679688 ) 96900.000000 , 9.690000 , 0.000000 ( 96909.687500 ) 97000.000000 , 9.700000 , 0.000000 ( 97009.703125 ) 97100.000000 , 9.710000 , 0.000000 ( 97109.710938 ) 97200.000000 , 9.720000 , 0.000000 ( 97209.718750 ) 97300.000000 , 9.730000 , 0.000000 ( 97309.726563 ) 97400.000000 , 9.740000 , 0.000000 ( 97409.742188 ) 97500.000000 , 9.750000 , 0.000000 ( 97509.750000 ) 97600.000000 , 9.760000 , 0.000000 ( 97609.757813 ) 97700.000000 , 9.770000 , 0.000000 ( 97709.773438 ) 97800.000000 , 9.780000 , 0.000000 ( 97809.781250 ) 97900.000000 , 9.790000 , 0.000000 ( 97909.789063 ) 98000.000000 , 9.800000 , 0.000000 ( 98009.796875 ) 98100.000000 , 9.810000 , 0.000000 ( 98109.812500 ) 98200.000000 , 9.820000 , 0.000000 ( 98209.820313 ) 98300.000000 , 9.830000 , 0.000000 ( 98309.828125 ) 98400.000000 , 9.840000 , 0.000000 ( 98409.843750 ) 98500.000000 , 9.850000 , 0.000000 ( 98509.851563 ) 98600.000000 , 9.860000 , 0.000000 ( 98609.859375 ) 98700.000000 , 9.870000 , 0.000000 ( 98709.867188 ) 98800.000000 , 9.880000 , 0.000000 ( 98809.882813 ) 98900.000000 , 9.890000 , 0.000000 ( 98909.890625 ) 99000.000000 , 9.900000 , 0.000000 ( 99009.898438 ) 99100.000000 , 9.910000 , 0.000000 ( 99109.906250 ) 99200.000000 , 9.920000 , 0.000000 ( 99209.921875 ) 99300.000000 , 9.930000 , 0.000000 ( 99309.929688 ) 99400.000000 , 9.940000 , 0.000000 ( 99409.937500 ) 99500.000000 , 9.950000 , 0.000000 ( 99509.953125 ) 99600.000000 , 9.960000 , 0.000000 ( 99609.960938 ) 99700.000000 , 9.970000 , 0.000000 ( 99709.968750 ) 99800.000000 , 9.980000 , 0.000000 ( 99809.976563 ) 99900.000000 , 9.990000 , 0.000000 ( 99909.992188 ) 続行するには何かキーを押してください . . .
当然ですが、プログラム中でオフセットを設定し、更に配列の最後まで計算させていないので、
計算していない場所は、初期化した時のまま0です。
計算結果の表示は、長くなるので、始めの50個と最後の50個だけです。

const size_t global_work_offset[]={30};
を
const size_t global_work_offset[]={0};
と書き換えると、オフセットが無くなります。
また、データ数1000に対して、
const size_t global_work_size[]  ={300};//タスクの総数
としているので、GPUは途中までしか計算していませんが、
const size_t global_work_size[]  ={1000};//タスクの総数
とすればGPUは最後まで計算してくれます。


Webで表示させる場合、どうしてもタブが上手く表示できないので、スペースキーに置換していたり、
色分けがされていなかったりと、
コードが見ずらいと思うので、一応プロジェクトを置いておきます。

ダウンロード先
MD5 | 372f9df2418c062753b2e06b0b4f5346 | 9.7 MB | OpenCL_Ver0.02_01
http://ux.getuploader.com/ADMIS/download/40/OpenCL_Ver0.02_01.zip


[2014_06_03]追記
他のプログラムに組み込み易いように、
もう少し処理を関数にまとめて、
main()関数を軽量化したバージョン。
(処理内容は同じ。)

MD5 | 7ad613859112e16ba7e37da1a9994d09 | 9.5 MB | OpenCLによるGPGPU
http://ux.getuploader.com/ADMIS/download/46/OpenCL_Ver0.02_02.zip