2010年9月1日
#
プロダクト受け入れテストと次版の設計のため
急遽デッチ上げられたプロジェクトに飛び込むことになり、
机のお引越ししてました。
東京タワーすぐそば一等地26階角部屋、
こんなの↓を背にシゴトしてます。
右隅にスカイツリーがほんのり。
そゆわけで、先週日曜は兎さんとやきそば食って
悟空茶荘で茶ぁしばきながら愛と平和を熱く語るはずもなく、
GPGPUがどーしたこーしたなんつーおよそ茶席にはそぐわぬ
ネタで盛り上がったわけですが。
兎さん曰く：「ATI Streamてばホスト-デバイス間のデータ転送が速くねぇ」
...そぉなの?
GPGPUはホスト(マザボ)にぶっ挿したデバイス(グラボ)に対し
あれやこれやと命令投げて処理を行います。
で、ホストはデバイスの持ってるメモリに直接アクセスできません。
なので必要に応じてホスト→デバイス,ホスト←デバイスなメモリ転送
コマンドを発行し、その完了を待たにゃならんです。そのスピードがイマイチとのこと。
おうち帰ってやってみた。1024x1024個のfloat(=4MB)領域に対し1000回の
書き込み(ホスト→デバイス)と読み込み(ホスト←デバイス)にかかる時間を計ってみる。
#include <iostream>
#include <vector>
#include <string>
#include <utility>
#include <cassert>
#include <Windows.h>
#define __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>
using namespace std;
int main() try {
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
assert( !platforms.empty() );
cl_context_properties properties[] =
{ CL_CONTEXT_PLATFORM, (cl_context_properties)(platforms[0])(), 0};
cl::Context context(CL_DEVICE_TYPE_GPU, properties);
std::vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
assert( !devices.empty() );
const size_t N = 1024*1024;
vector<float> host_buffer(N, 0.0f);
cl::Buffer buffer(context, CL_MEM_READ_WRITE, N*sizeof(float));
cl::CommandQueue queue(context, devices[0]);
DWORD t0, t1;
// かきかき
t0 = GetTickCount();
for ( int i = 0; i < 1000; ++i ) {
queue.enqueueWriteBuffer(buffer, CL_TRUE, 0, N*sizeof(float), &host_buffer[0]);
}
t1 = GetTickCount();
cout << "Write 1000 times: " << (t1-t0) << "[ms]\n";
// よみよみ
t0 = GetTickCount();
for ( int i = 0; i < 1000; ++i ) {
queue.enqueueReadBuffer(buffer,CL_TRUE, 0, N*sizeof(float), &host_buffer[0]);
}
t1 = GetTickCount();
cout << "Read 1000 times: " << (t1-t0) << "[ms]\n";
} catch ( const cl::Error& err ) {
cerr << "OpenCL error " << err.err() << " :" << err.what() << endl;
}
ATI Stream SDK 2.2 / ATI Radeon HD5450 DDR2-512MB での結果:
Write 1000 times: 11766[ms]
Read 1000 times: 5313[ms]
ほほー、書き込みは読み込みの倍以上の時間がかかるですね。
ホスト-デバイス間のメモリ転送(特にホスト→デバイス)はパフォーマンスを落とすのか。
2010年8月29日
#
ひさしぶりの茶ネタ。
今年の夏休みも終わったです。明日からおしごと、行きたくないよぅ...
わんくま同盟横浜勉強会、たくさんのオーディエンスに集まっていただきました。
懇親会も盛況でした、ホントありがとうございます。
勉強会当日夜行バスで岩手からすっ飛んで来てくれた兎さんに
何のお相手もできなかったので、今日のお昼はおもてなし。
つーてもいつものごとく茶ぁ買って茶ぁしばいただけですが。
昼飯はネタ度の高い梅蘭やきそばなどを。
兎さんの茶買い出しツアーのついでに僕もいくつか手に入れて、
その中の一つ：福建青茶「白牙奇蘭」
岩茶に仕立てたものもありますが、こいつはゆるく丸めたタイプ。
福建の茶らしく、鉄観音あるいは黄金桂系の香味です。
ぁー、ひっさしぶりに真面目に茶淹れた。ｳﾏｰ
2010年8月21日
#
今日からεπιは夏休みです。
ウチとこはまとめて一週間、なつやすみを好きに選べるです。
ここんとこ何年も、Tech・Edに合わせてますけどね。
なつやすみの間、どこぞに旅に出る予定もなく、てろてろとコード書いて遊んでんじゃないかな。
※ 仕事のコードと遊びのコードは別モノです。仕事のコードはやんちゃが許されないのでつまんなーい。
水曜からTech・Ed 始まるんでコーディング三昧、てほどではないけども、
数日間は手を染め始めたOpenCLをこねこねする予定。
いつか勉強会でお披露目したいとこなんだが、デモできないのが辛いとこ。
マンデルブロ集合を瞬間で描き切るとこ、お見せしたいんですけどねー。
2010年8月18日
#
ぼくらの夏祭り: Microsoft Tech・Ed 2010 まであと一週間と迫りました。旅支度は
恙無く進んでおりますでしょうか。
わんくま同盟の見飽きた見慣れた連中はBOFとExpert Talkあたりに出没します。
僕はっちゅーと今年は残念ながらBOFの出番ナシ。
んだけどおそらくはBOF会場に顔出して粛々と進行するセッションをひっかき回す気マンマンでございますとも。
僕の出番はExpert Talkっす。よーするに「よろず相談承り処」ですね。MSのスタッフとMVPが皆さんの質問/相談にお答えしようって企画です。ここもえらくわんくま色のこゆーい一角となってます。僕は初日(8/25)15:20-16:55と三日目(8/27)13:45-1520に座ってろってお達し。
んなわけで御来場の皆さんよろしくどーぞ。
んでもってついでに土曜日(8/28)はわんくま同盟横浜勉強会「C++Day」です。今回もこゆーいネタを用意しましたですよー
2010年8月16日
#
んーむ、参ったね。
OpenCL 1.1 を試すべく 最新ドライバ Catalyst 10.7 を突っ込んで、
しばらくは快調に動いてくれてたんだけども、なにが哀しいのか
困ったトラブル。使用中突然画面が固まったり暗転したり。
ドライバ更新の際はinstall manager CCC をun-install するが無難
とか風のうわさで聞いてやってみたけど解消せず。
safe-modeでドライバ抜いてVGA互換モードでなら問題ないので、
ハードがイカれたとは考えにくい。
ショボいHD5450に差し替えて様子をみることにするつもりだけど...
OpenCLさっくり動いちゃったのでゴキゲンだったのになー。しょんぼりさんです。
もうちょい上位のになると補助給電せにゃならんので電源取り換えだし。
てか今の電源がグラボ電源持ち合わせてないからコイツにしたのにのに。
なんやわけわからんトラブルです。
なんか情報あったら教えてくださいまし > ALL
OpenCLではソースコードを文字列で与え、実行時にコンパイルしています。
こうすることで、Platformが提供する各種デバイスに応じたProgramを用意できるわけですが、いかんせんコンパイル時間がMOTTAINAI。デバイスを決め打ちにしていいならコンパイル済のBinary Kernelを直接ロードするんが手っ取り早いす。
で、Binary Kernelのつくりかた。
KB115 - ATI Stream SDK v2.2 Support for Binary OpenCL™ Kernels
のキモんとこを解説。
0. 前準備
上記KB115ページに clbinary.zip
があります。こいつを落っことして解凍し、clbinarygen.c, clbinaryuse.c を用意します。
1. clbinarygen, clbinaryuse をコンパイルし、exeを作ります。
vc10ならば:
: build.bat
set INCDIR="%ATISTREAMSDKROOT%include"
set LIBDIR="%ATISTREAMSDKROOT%lib\x86"
cl -EHsc -I%INCDIR% %1 %2 %3 %4 %5 OpenCL.lib -link -libpath:%LIBDIR%
なんてなバッチからコンパイルすりゃえぇべ。
2. カーネル・コードをコンパイルし、Kernel Binaryを作ります。
カーネル・コード ほげほげ.cl を kernel.cl にコピーし、clbinarygen を起動すると:
>clbinarygen
Reading in program from kernel.cl
Number of platforms found: 1
Found AMD platform
Number of devices found: 10
DEVICE[0]: CL_DEVICE_TYPE_CPU, Name=AMD Athlon(tm) 64 X2 Dual Core Processor 4400+
DEVICE[1]: CL_DEVICE_TYPE_GPU, Name=Redwood
DEVICE[2]: CL_DEVICE_TYPE_GPU, Name=ATI RV770
DEVICE[3]: CL_DEVICE_TYPE_GPU, Name=ATI RV770
DEVICE[4]: CL_DEVICE_TYPE_GPU, Name=ATI RV710
DEVICE[5]: CL_DEVICE_TYPE_GPU, Name=ATI RV730
DEVICE[6]: CL_DEVICE_TYPE_GPU, Name=Cypress
DEVICE[7]: CL_DEVICE_TYPE_GPU, Name=Juniper
DEVICE[8]: CL_DEVICE_TYPE_GPU, Name=Redwood
DEVICE[9]: CL_DEVICE_TYPE_GPU, Name=Cedar
Writing out binary kernel to kernel.bin.0
Writing out binary kernel to kernel.bin.1
Writing out binary kernel to kernel.bin.2
Writing out binary kernel to kernel.bin.3
Writing out binary kernel to kernel.bin.4
Writing out binary kernel to kernel.bin.5
Writing out binary kernel to kernel.bin.6
Writing out binary kernel to kernel.bin.7
Writing out binary kernel to kernel.bin.8
Writing out binary kernel to kernel.bin.9
こんなん出ました。各Device対応に kernel.bin.xxx が生成されます。僕のRADEON HD5670 のDevice名は Redwood だから kernel.bin.1 か kernel.bin.8 を使えばいいみたいよ。
3. できあがったBinary Kernelはそのまんま使えるみたいだけど、そこそこデカいので不要なデータをさっぴいて小さくします。
Binary Kernelの要らんとこをそぎ落として小さくするにはLinux由来のツール: readelf, objcopy が必要となります。僕はWindowsのヒトなのでCygwinをインストールしておきます。そんときDevelカテゴリのbinutilsパッケージをお忘れなく。この中にreadelf,objcopyがありますでの。
>c:\cygwin\bin\readelf -aW kernel.bin.1 | c:\cygwin\bin\grep Machine
Machine: <unknown>: 0x3f3
こんなん出ました。ケツにある１６進: 0x3f3 が大事。
んでは要らん部分をそぎ落とします。
32bit:
objcopy -I elf32-i386 -O elf32-i386 -R .source -R .livmir -R .amdil --alt-machine-code=<さっきの１６進コード> <Kernel Binary名> <できあがり>
64bit:
objcopy -I elf64-x86-64 -O elf64-x86-64 -R .source -R .livmir -R .amdil --alt-machine-code=<さっきの１６進コード> <Kernel Binary名> <できあがり>
僕んトコだと:
>c:\cygwin\bin\objcopy -I elf32-i386 -O elf32-i386 -R .source -R .livmir -R .amdil --alt-machine-code=0x3f3 kernel.bin.1 stripped_kernel.bin
/usr/bin/objcopy: this target does not support 1011 alternative machine codes
/usr/bin/objcopy: treating that number as an absolute e_machine value instead
なんかゆってるけどひとまず気にしないｗｗ
4. できあがったstripped_kernel.binが使えることを確認します。
>clbinaryuse
Number of platforms found: 1
Found AMD platform
Number of devices found: 2
DEVICE[0]: CL_DEVICE_TYPE_CPU, Name=AMD Athlon(tm) 64 X2 Dual Core Processor 4400+
DEVICE[1]: CL_DEVICE_TYPE_GPU, Name=Redwood
Which device do you want to use? [0-1]?
1
Enter filename of corresponding binary kernel file
stripped_kernel.bin
Reading in binary kernel from stripped_kernel.bin
*** REPLACE THIS WITH ACTUAL WORK ***
↑コレが見えたらできあがり。
このバイナリ・ファイルを読み込むには:
cl::Context context;
vector<cl::Device> devices;
...
vector<char> binary;
{
ifstream stream("stripped_kernel.bin", ios::binary);
if ( !stream.is_open() ) { return 0; }
istreambuf_iterator<char> first(stream);
istreambuf_iterator<char> last;
copy(first, last, back_inserter(binary));
}
cl::Program::Binaries binaries;
binaries.push_back(make_pair(static_cast<void*>(&binary[0]),binary.size()));
cl::Program program(context,devices,binaries);
とかなんとか。
2010年8月14日
#
前回の奇遇転置ソートではトホホな結果となりました。
あまりに悔しいのでマンデルブロ集合に挑戦。
マンデルブロ集合は
int getCalcNum(const float x, const float y, const int max_calc) {
int count = 0;
float prevX = 0.0f;
float prevY = 0.0f;
for (int i = 0; i < max_calc + 1; i++) {
count = i + 1;
float newX = (prevX * prevX) - (prevY * prevY) + x;
float newY = (2.0f * prevX * prevY) + y;
if (((newX * newX) + (newY * newY)) > 4.0f) {
// 発散する
break;
} else {
prevX = newX;
prevY = newY;
}
}
return count;
}
こんな計算を各ピクセルに対して行います。
Kernelに計算範囲と結果格納領域を与えて起動すれば、
あとはホストとのやり取りなしに全領域を計算して返ってくるよに仕向けました。
↑の計算を512x512点、max_calc=512 でぶん回してみた。
結果:
single-thread@CPU : 180[ms] （AMD Athlon64 X2 4400+)
OpenCL : 15[ms] (ATi RADEON HD5670)
こぉでなくちゃだわ。
時間計測にはおちゃらかにGetTickCount()使こうたんで精度はめちゃﾃｷﾄｰですが、
こんくらいのパフォーマンスは叩き出せるですね。GPGPUはやればできる子。
2010年8月13日
#
ここ数日でOpenCLについて学んだこと。
- 各ベンダはOpenCLのPlatformを提供する。
- Platformはひとつ以上のDeviceを持つ。
- 特定のDeviceで実行する環境をContextという。
- Deviceに与える(OpenCL Cで書かれた)Programはひとつ以上のエントリ(関数の入り口)を持つ。
- ひとつのエントリとそれに与える引数のセットをKernelという。
- Device内のBuffer(メモリ領域)への読み書き、
そしてKernelによる実行などをCommandQueueを介して行う。
これらを踏まえてこしらえた習作、お題は「並列処理が向かないかもしれない例」の奇遇転置ソート:
1 7 6 5 8 2 4 3 があったとき、
[1] ふたつずつのペアをつくる。1 7 6 5 8 2 4 3
[2] それぞれのペアについて必要なら入れ替え 1 7 5 6 2 8 3 4
[3] ひとつずらしてペアをつくる 1 7 5 6 2 8 3 4
[4] それぞれのペアについて必要なら入れ替え1 5 7 2 6 3 8 4
[5] [1]～[4] をくりかえし。
[2],[4] の際、ペアの入れ替えは互いに独立だから並行処理できんぢゃん、と。
なんたってGPUですからね、コアはたっぷりあります。
目下愛用のRADEON HD5670、シェーダ・プロセッサを400基載せてます。
#include <iostream> // cout, endl
#include <array> // array
#include <vector> // vector
#include <map> // map
#include <string> // string
#include <numeric> // iota
#include <algorithm> // copy, for_each
#include <iterator> // back_inserter
#include <utility> // make_pair
#include <cassert> // assert
#include <Windows.h> // GetTickCount
// error時は例外処理に一任する
#define __CL_ENABLE_EXCEPTIONS
#include <CL/cl.hpp>
using namespace std;
int main() {
try {
// Platformを列挙し、ひとつ目を使う
vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
assert( !platforms.empty() );
cl::Platform platform = platforms[0];
cout << "Platform: " << platform.getInfo<CL_PLATFORM_NAME>() << '\t';
// 最初に見つかったPlatformからデバイスを探し、GPU-Contextを生成
cl_context_properties cprops[3] = { CL_CONTEXT_PLATFORM,(cl_context_properties)platform(), 0};
cl::Context context(CL_DEVICE_TYPE_GPU,cprops);
// Context中のDeviceを列挙し、ひとつ目を使う
vector<cl::Device> devices = context.getInfo<CL_CONTEXT_DEVICES>();
assert( !devices.empty() );
cl::Device device = devices[0];
cout << "Device: " << device.getInfo<CL_DEVICE_NAME>() << endl;
// DeviceにCommandQueueを用意
cl::CommandQueue queue(context,device);
// Programを準備 : 奇遇ソート
static string source_strs[] = {
"__kernel void swap_even( \n"
" __global float* data, \n"
" __global int* swapped) \n"
"{ \n"
" int i = get_global_id(0); \n"
" int j = i + 1; \n"
" if ( !(i & 1) ) { \n"
" if ( data[i] > data[j] ) { \n"
" float tmp = data[i]; \n"
" data[i] = data[j]; \n"
" data[j] = tmp; \n"
" *swapped = 1; \n"
" } \n"
" } \n"
"} \n"
,
"__kernel void swap_odd( \n"
" __global float* data, \n"
" __global int* swapped) \n"
"{ \n"
" int i = get_global_id(0); \n"
" int j = i + 1; \n"
" if ( (i & 1) ) { \n"
" if ( data[i] > data[j] ) { \n"
" float tmp = data[i]; \n"
" data[i] = data[j]; \n"
" data[j] = tmp; \n"
" *swapped = 1; \n"
" } \n"
" } \n"
"} \n"
,
""
};
cl::Program::Sources sources;
for ( string* str = source_strs; !str->empty(); ++str ) {
sources.push_back(make_pair(str->c_str(), str->size()));
}
// Programのビルド(失敗時にはログを出力)
cl::Program program(context, sources);
try {
program.build(devices);
} catch ( cl::Error& ) {
cerr << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << endl;
return -1;
}
// ProgramからKernelを生成
map<string,cl::Kernel> kernels;
{
vector<cl::Kernel> knls;
program.createKernels(&knls);
// 関数名とKernelとの対応表をつくる
for_each(knls.begin(),knls.end(),
[&](cl::Kernel k) { kernels[k.getInfo<CL_KERNEL_FUNCTION_NAME>()] = k; });
}
// ホストに要素数Nの配列を用意
const int N = 10000;
array<float,N> data;
int swapped;
// data[] = { N-1, N-2, ... 2, 1 }
iota(data.begin(), data.end(), 1.0f);
reverse(data.begin(),data.end());
// ふつーにやればこんなの。
DWORD t = GetTickCount();
do {
swapped = 0;
for ( int i = 0; i < N-1; i += 2 ) {
if ( data[i] > data[i+1] ) {
float tmp = data[i]; data[i] = data[i+1]; data[i+1] = tmp;
swapped = 1;
}
}
for ( int i = 1; i < N-1; i += 2 ) {
if ( data[i] > data[i+1] ) {
float tmp = data[i]; data[i] = data[i+1]; data[i+1] = tmp;
swapped = 1;
}
}
} while ( swapped );
t = GetTickCount() - t;
assert( is_sorted(data.begin(), data.end()) );
if ( is_sorted(data.begin(),data.end()) ) {
cout << "CPU DONE. " << t << "[ms]" << endl;
}
reverse(data.begin(),data.end());
// Device側にバッファを確保
cl::Buffer buffer_data(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
sizeof(data) , &data[0]);
cl::Buffer buffer_swapped(context, CL_MEM_WRITE_ONLY,
sizeof(swapped), nullptr);
// 2つのKernelに引数を設定
cl::Kernel kernel_0 = kernels["swap_even"];
kernel_0.setArg(0,buffer_data);
kernel_0.setArg(1,buffer_swapped);
cl::Kernel kernel_1 = kernels["swap_odd"];
kernel_1.setArg(0,buffer_data);
kernel_1.setArg(1,buffer_swapped);
t = GetTickCount();
do {
// swappedを0(false)に
swapped = 0;
queue.enqueueWriteBuffer(buffer_swapped,CL_TRUE, 0, sizeof(swapped), &swapped);
// 偶数
queue.enqueueNDRangeKernel(kernel_0, cl::NullRange, cl::NDRange(N-1), cl::NullRange);
queue.enqueueBarrier();
// 奇数
queue.enqueueNDRangeKernel(kernel_1, cl::NullRange, cl::NDRange(N-1), cl::NullRange);
queue.enqueueBarrier();
// swappedを取り出し
queue.enqueueReadBuffer(buffer_swapped, CL_TRUE, 0, sizeof(swapped), &swapped);
} while ( swapped );
t = GetTickCount() - t;
// 結果を取り出して検証
queue.enqueueReadBuffer(buffer_data, true, 0, sizeof(data), &data[0]);
assert( is_sorted(data.begin(), data.end()) );
if ( is_sorted(data.begin(), data.end()) ) {
cout << "GPU DONE. " << t << "[ms] " << endl;
}
} catch ( cl::Error& ex ) {
cerr << ex.err() << ':' << ex.what() << endl;
return -1;
}
}
実行結果:
Platform: ATI Stream Device: Redwood
CPU DONE. 156[ms]
GPU DONE. 2886[ms]
あ、あるえぇぇ...? ぜんっぜん速くねぇ。
赤いトコがGPU内でチョー並列に動いてはくれるんだけども、
それを呼んでる青いトコはホスト側にあります。
青いトコは並列でもなんでもなく、赤いトコがいっくら速くても
青いトコがそれを相殺してるように思えます。
GPGPUでかっとびアプリを書くにはまだまだ修行が足らんなーと。
まぁ二、三日でここまで書けるよになったんだから僕の感覚もそんなに鈍っちゃいないぜ、と。
2010年8月12日
#
あらま早々にSDKのupgrade.
ATI Stream Software Development Kit (SDK) v2.2
大きな変更点(What's New? より抜粋):
- OpenCL 1.1 サポート
- Visual Studio 2010 に対応 (残念ながらExpress.edはだめポ)
- double型 サポート (#pragma OPENCL EXTENSION cl_amd_fp64 : enable がおまぢない)
- debug用に printf できる(#pragma OPENCL EXTENSION cl_amd_printf : enable)
などなど。
早速落っことしてサンプルを試してるとこっす。VS10用のsolutionがついてました。
# 対応ドライバは Catalyst 10.7 です。
2010年8月10日
#
CodeZineです、はい。
shared_ptrによるObserverパターンの実装
お馴染みObserverパターンです。要はCallbackすよ。
Observerの歴史は古く、織田信長が長引く戦に疲弊し、矢の飛び交う中
「わしはちょっと寝る。なにかあったら起こせ」
と臣下に伝え馬上で熟睡したとの記録が残っており、
馬上で眠る信長公の姿を模した木彫りの馬: 御武座馬(おぶざば)
は郷土玩具として有名ですね。(うっそぴょーん)
Observerパターンに起こる問題のひとつが、
「起こそうとしたお館様が冷たくなってたらどーすんの?」
.NET/Javaでは強参照、つまり「死なさへんでぇ」と。
C++では参照カウントを使って死なさんこともできますし、
弱参照を使えば死んでもだいぢょぶな実装もできるよ。
ってなことを書いてます。コードたっぷりです。
ぢつは、このお話もう一本のルートがありましてですね。
.NETではdelegate/eventってゆーナイスなからくりが
Observerパターンの利用を自由なものにしてくれました。
delegateのおかげで特定のクラス/インタフェースから継承/実装
しなければ信長になれないってゆー縛りから解放されます。
さて、C++0Xでは新たにlambdaが導入されました。
このlambdaを使ってdelegate/eventの真似ごとができるはず。
どんな手管でdelegate/eventっポいものを作るか...知りたくね?
わんくま横浜勉強会にいらっしゃいませ(ﾎﾎﾎ)。
マンガ本とか魔導書とか怪しげなカードとかもありますよー
なんとなくOpenCLなるものを の続き。
そゆわけで、入門書を読みつつ第三章「OpenCLプログラム作成手順」に従って
サンプル・コードを打ち込んでたわけです。
版社のサポート・ページから掲載コードぜんぶ拾ってこれるんだけども、
最初はちまちま手打ちして開発スタイルを腕に覚えさすですよ。
サンプル・コードはなんとも素直なCなので難しくはないんだけども...長い。
二つの float 配列を足す: out[i] = in1[i] + in2[i] をやらすだけのことに400行。
毎回これじゃかなわんなー、なんぞ賢いやり方ないものかとATIやOpenCLの
本家クロノスにある資料を眺めてたですよ。
...みつけましたぜダンナ、クロノスのOpenCLドキュメント集の中に
OpenCL 1.1 C++ Bindings Specification なんちゅーのがあるやないのん。
んむ、どうやらCで提供されたOpenCL-APIにクラスの薄皮を被せたwrapperとミタ。
そんじゃってんで ATI Stream SDK の include-dir. 覗いたらアナタ、ちゃんと入ってるぢゃあァりませんか♪
こいつを使って書き換えてみたらまぁ、なんということでしょう。400行が80行にっ!
おちゃらか超並列コンピューティングなら OpenCL C++binding でキマリだぜぃ!
# ここらへん、.NETはどんなアプローチなのかしら。
# DirectCompute とやらの .NET-interface を提供すんのかしら。
2010年8月8日
#
お散歩がてら横浜近くの本屋さん:有隣堂＠関内をぶらついてきました。
# ココは文具の品ぞろえもよろしく、お気に入りです。
なんとはなしに買ってきた「OpenCL入門」←衝動買いっす。
いやね、お友達から「CUDAやんね?」ってお誘いを受け、グラボ(GeForce9800GTS)まで借りてたですよ。
ところがそのグラボ、こんがり焼けちゃいまして。
かわりに買ったのがRADEON HD5670積んだやつ。能力的にはとんとんかな体感では。
結局CUDAに触ることなくNVIDIAからATIに転んだわけで、
CUDAのお誘いに乗ることができず少なからず負い目を感じておったわけっす。
その罪滅ぼし、てーことでもないけどOpenCLならCUDA/ATI-streamを包含すっからちぃとは言い訳もたつんでないかと。
それにGPGPU絡みはCUDAが先行してたんでATIユーザには情報がまだ少ないポ。
僕が触ればちっとはお役に立てんでないかと。
まずは OpenCL 開発前準備 (ATIのグラボ持ってる人向け)
1. ATI Stream のSDK入手とinstall
AMDのサイト: ATI Stream Software Development Kit (SDK) v2.1 から SDKを拾ってくる。
各種環境に合わせたSDKがあるんでﾃｹﾄｰに選んでくだち。
僕は Vista/Win7版の32bit/64bitの両方落としといたです。
開発環境はVS2008となってますがVS2010でも問題なさげ。
ATI Streamで遊べるGPUリストもこのページに載ってるんでざっくし確認するがよいー(RADEON HD4xxx/5xxxなら大抵OKみたいです)。
インストールはインストーラが上がってくれるんで[次へ]ボタンを突っついただけです。
2. コンパイル・オプションを調べる
SDKに同梱されてたサンプル・ソリューションを開き、いくつかbuild/実行して動くのを確認。
しかるのちやおらプロジェクト・プロパティこじあけてコンパイル・オプションを眺める。
いやさ、いちいちVS立ち上げるよかテキストエディタとコマンドライン・コンパイラでちまちま
いぢりたいこともあるんでね。
...どうやらSDKインストールした時点で環境変数 ATISTREAMSDKROOT が定義されたみたい。
ここらへんの情報を基に build.bat をこしらえた:
set INCDIR="%ATISTREAMSDKROOT%include"
set LIBDIR="%ATISTREAMSDKROOT%lib\x86"
cl -EHsc -I%INCDIR% %1 %2 %3 %4 %5 OpenCL.lib -link -libpath:%LIBDIR%
こんなのをテキトーに掘ったディレクトリ d:\work\OpenCL に置く。
3. おためし
買ってきた「OpenCL入門」のサポートページからサンプルをひろって解凍。いっちゃん最初のサンプル main.cpp をさっきのディレクトリにコピっていきなりbuildを試みる。
d:\Work\OpenCL>build main.cpp
d:\Work\OpenCL>set INCDIR="C:\Program Files (x86)\ATI Stream\include"
d:\Work\OpenCL>set LIBDIR="C:\Program Files (x86)\ATI Stream\lib\x86"
d:\Work\OpenCL>cl -EHsc -I"C:\Program Files (x86)\ATI Stream\include" main.cpp
OpenCL.lib -link -libpath:"C:\Program Files (x86)\ATI Stream\lib\x86"
Microsoft(R) 32-bit C/C++ Optimizing Compiler Version 16.00.30319.01 for 80x86
Copyright (C) Microsoft Corporation. All rights reserved.
main.cpp
Microsoft (R) Incremental Linker Version 10.00.30319.01
Copyright (C) Microsoft Corporation. All rights reserved.
/out:main.exe
"-libpath:C:\Program Files (x86)\ATI Stream\lib\x86"
main.obj
OpenCL.lib
なんかできたよ。実行してみっか:
d:\Work\OpenCL>main
Number of platform(s) : 1
Platform profile : FULL_PROFILE
Platform version : OpenCL 1.0 ATI-Stream-v2.1 (145)
Platform name : ATI Stream
Platform vendor : Advanced Micro Devices, Inc.
Platform extensions: cl_khr_icd
...さっくりイケちゃったみたい。出だし好調っすねー。
2010年8月7日
#
ボーズにおねだりされまして。
「ドミニオンのサプライから10枚ﾃｹﾄｰに選んでくれるアプリ作ってぇ」
「作っちゃるよー、んじゃ週末にちゃっちゃとね」
なんつー安請け合い。ま、どってことないんじゃね(ほじほじ)
元ネタとなるカードデータはツブシが効くようにXMLで用意しました:
<?xml version='1.0' encoding='shift_jis' ?>
<dominion>
<card>
<key>0.01</key>
<name>へそくり</name>
<series>プロモ</series>
<cost>5</cost>
<portion_cost>0</portion_cost>
<action>false</action>
<treasure>true</treasure>
<point>false</point>
<attack>false</attack>
<reaction>false</reaction>
<continuous>false</continuous>
</card>
<card>
<key>1.01</key>
<name>市場</name>
<series>無印</series>
<cost>5</cost>
<portion_cost>0</portion_cost>
<action>true</action>
<treasure>false</treasure>
<point>false</point>
<attack>false</attack>
<reaction>false</reaction>
<continuous>false</continuous>
</card>
...
</dominion>
.NETならSystem.Xml.Linqで
// dominion.xmlから<card>要素を列挙して
foreach ( XElement element in XElement.Load("dominion.xml").Elements("card") ) {
どーのこーの
}
すりゃえぇわけだが、ボーズのマシンがショボくてメモリかつかつで.NETなんつー大それたもんインスコしたくないとゆー。うんじゃnativeで書くかと。
ってわけでひさしぶりにApache Xerces-Cとかひっぱりだしてきたりする。
まずは手習いがてら<name>でも列挙してみっか：
#include <iostream>
#include <locale>
#include <memory>
#include <xercesc/parsers/XercesDOMParser.hpp>
#include <xercesc/dom/DOM.hpp>
#include <xercesc/sax/HandlerBase.hpp>
#include <xercesc/util/XMLString.hpp>
#include <xercesc/util/PlatformUtils.hpp>
using namespace std;
using namespace xercesc;
int run(const wchar_t* target) {
unique_ptr<XercesDOMParser> parser(new XercesDOMParser());
unique_ptr<ErrorHandler> errHandler(static_cast<ErrorHandler*>(new HandlerBase()));
parser->setErrorHandler(errHandler.get());
try {
// パースして
parser->parse(target);
// ドキュメントのルートエレメント配下の<name>要素を抜き出して
DOMNodeList* nodes = parser->getDocument()->getDocumentElement()->getElementsByTagName(L"name");
// テキスト・コンテントをプリントしるー
for ( XMLSize_t i = 0; i < nodes->getLength(); ++i ) {
wcout << dynamic_cast<DOMElement*>(nodes->item(i))->getTextContent() << endl;
}
} catch ( ... ) {
wcout << L"Something wrong" << endl;
return -1;
}
return 0;
}
int main() {
wcout.imbue(locale("japanese"));
try {
XMLPlatformUtils::Initialize();
} catch ( const XMLException& ex ) {
wcout << ex.getMessage() << endl;
return -1;
}
int result = run(L"dominion.xml");
XMLPlatformUtils::Terminate();
return result;
}
実行結果:
へそくり
市場
改築
鍛冶屋
金貸し
木こり
議事堂
玉座の間
研究所
鉱山
工房
宰相
....
VS2010はnative-C++でのｲﾝﾃﾘｾﾝｽがひじょーにキモチ良く動いてくれるんで快適です。
ちょいと昔はマニュアル首っぴきで苦労したですが。
2010年8月6日
#
TBBの並列アルゴリズム parallel_なんちゃら の中から使用頻度の高そげなparallel_reduce をご紹介。
まぁキホン parallel_for と同じく繰り返し(loop)の並列化なんだけども parallel_reduce には繰り返しの合間に reduction(リダクション:還元/縮約/要は"おまとめ")が割り込みます。
ふたつのベクトル x[0..N-1] と y[0..N-1] の内積を求めてみんとす。内積ちゅーのは:
x・y = ∑x[i]*y[i] (i = 0..N-1)
ですな。こいつはインデクスを 0..M-1 と M..N-1 の二つに分けて
∑x[i]*y[i] (i = 0..M-1) と
∑x[i]*y[i] (i = M..N-1) との
和を求めても同じ結果。いくつに分割してもかまわんから並列処理にはうってつけ。
ただし、それぞれの部分和を積算(おまとめ)するとこ(ココがreduction)では各スレッドが勝手にやらんようにガードせんならんです。
parallel_reduceにはloop範囲と初期値、処理本体およびreduction処理を与えます。
#include <iostream>
#include <array>
#include <numeric>
#include <tbb/tbb.h>
using namespace std;
// 作為的に遅い掛け算
double mult(double x, double y) {
const int N = 100000;
double sum = 0.0;
for ( int i = 0; i < N; ++i ) {
sum += x*y;
}
return sum / (double)N;
};
// フツーに内積
tbb::tick_count::interval_t
serial_inner_product(double x[], double y[], double& z, int n) {
tbb::tick_count t = tbb::tick_count::now();
double result = 0.0;
for ( int i = 0; i < n; ++i ) {
result += mult(*x,*y);
++x; ++y;
}
z = result;
return tbb::tick_count::now() - t;
}
// 並列に内積
tbb::tick_count::interval_t
parallel_inner_product(double x[], double y[], double& z, int n) {
tbb::tick_count t = tbb::tick_count::now();
z = tbb::parallel_reduce( tbb::blocked_range<int>(0,n), 0.0, // [0,n) の範囲で、初期値0.0
// 切り分けられた範囲での内積を求め
[&](const tbb::blocked_range<int>& range, const double& value) -> double {
double result = value;
for ( int i = range.begin(); i != range.end(); ++i ) {
result += mult(x[i],y[i]);
}
return result;
},
// リダクション汁
[](const double& x, const double& y) { return x + y; });
return tbb::tick_count::now() - t;
}
int main() {
// 1*1 + 2*2 + 3*3 + ... + N*N を求めるよ
const int N = 100;
array<double,N> x;
array<double,N> y;
iota(x.begin(),x.end(),1.0);
iota(y.begin(),y.end(),1.0);
double sz, pz;
std::cout << "serial: " << serial_inner_product(&x[0],&y[0],sz,N).seconds() << "[sec.]\n";
std::cout << "parallel: " << parallel_inner_product(&x[0],&y[0],pz,N).seconds() << "[sec.]\n";
std::cout << sz << '/' << pz << std::endl;
}
実行結果(dual-core)
serial: 0.0807083[sec.]
parallel: 0.0465302[sec.]
338350/338350
おー、いぃ感じにcore数なりのスピードアップなりねー