OpenCLプログラミング#1

■ このスレッドは過去ログ倉庫に格納されています
1a36 ◆K0BqlCB3.k
垢版 |
2008/12/10(水) 15:38:25
さてついにOpenCLの仕様が公開されました。

http://www.khronos.org/opencl/

公式ページにはAPIのヘッダファイルが公開されており、
まだ実際に動かす事はできないもののプログラミングすることは可能となっています。
ということで、公開に先んじてプログラミングを始めてしまいましょう。
2011/08/04(木) 04:52:42.94
APP SDK 2.5 で BFI_INT が使われるようになったけど、
bitselect() 使うとだめで、 (x & 0xf0f0f0f0) | (y & 0x0f0f0f0f) の型だと OK みたい。
2011/08/04(木) 18:42:12.38
CAL ILコンパイラはもう誰もメンテできないんだろうな。BFI_INT を付け足すのがそんなに苦痛かね?
AMDはLLVMのラデバックエンドに挑戦しているようだけど、果たして新アーキに切り替わる前に
モノにできるんだろうか?
2011/08/04(木) 23:21:28.44
OpenCLってCellみたいなプロセッサが普及しないと面白くないよね。
GPU型の抽象化だけじゃ。
2011/08/05(金) 09:45:21.35
そのCellがコケたからなー
発想が速すぎた
今の技術で開発していれば、大人気だったかもしれないのに
2011/08/06(土) 00:40:12.65
i7であえてOpenCL
2011/08/06(土) 02:51:44.01
OpenCLのアプリを配布するときってAMDとnVidiaで別に実行ファイル作らなきゃいけないの？
2011/08/06(土) 03:08:35.23
何でんなこと思った。
2011/08/06(土) 03:57:16.00
今のところ、OpenCLなアプリケーションを"配布"するための良い方法は無い。(除 MacOS X)
2011/08/06(土) 07:16:37.83
*.cl カーネルを露出させるのがそんなに恥ずかしいか!

難読化するだけだったら gzip してバイナリに埋め込んどくのが楽かもね。
2011/08/07(日) 08:06:38.42
nVidiaから最新ドライバでOpenCL1.1対応とかメールきたんだけど、
ドライバとりに行ったらLinux用はまだじゃねーか。
ざけんなー。
2011/08/07(日) 08:08:49.28
>>480
恥ずかしいというよりノウハウの塊だから出したくない人はいるだろう。
2011/08/08(月) 05:22:20.10
Windows用OpenCL1.1ドライバを試してみたが、
どうも上手く動かないのでもう一寸待った方が良いと思う
2011/08/09(火) 01:37:46.04
>>428
門外不出のノウハウ＝世間一般には知られていない。
つまり各分野の研究者すら出し抜く大発見ってことだぞｗ
2011/08/11(木) 20:59:03.48
>>484
ノウハウが全て開示されているとでも思っているのか？
2011/08/11(木) 22:05:51.96
世間一般に知られていなければ十分じゃん
2011/08/12(金) 21:31:49.60
世間一般に知られないようにするためにカーネルを晒したくないわけだろ
2011/08/14(日) 01:34:57.52
>>473
CAL ILコンパイラ自体はいまだに新命令追加されているよ。
64bitINT除算とかはドキュメント化までされている。
ドキュメントにはなっていないけど、wave_idとかは
Appendix Bにはこっそり追加されて、実際テキストで
1OP命令として書けば使える。


でもBFI_INTは置いてけぼり
2011/08/14(日) 09:57:57.76
マスクが定数なら、iand,iand,ior で BFI_INT になるとおもうけど、
そうじゃないときになー。
2011/08/31(水) 17:10:54.28
初心者です。
Cで書かれた演算プログラムを高速化させたいと思っています。
この元のプログラムをホストプログラムとして、OpenCL用に書き換えて、高速化させたい演算部分を引っこ抜いてカーネルプログラムとして書き換えるというやり方でいいのでしょうか？

また、このようなCのプログラムをOpenCL用に書き換える作業を解説しているサイトや書籍はあるのでしょうか？

当方LINUXでのプログラミングをしております。
2011/08/31(水) 21:06:58.51
並列化はアルゴリズムをかなり選ぶよ。GPGPUなんて尚更。
まず普通にマルチスレッドで並列化してみたら?
2011/09/15(木) 10:02:24.21
opencl用の自動並列化コンパイラってある？
2011/09/16(金) 05:07:07.00
>>490
そのやりかたでいいさ。
2重ループを1重ループ2つにわけるとか。
あとはほとんどC言語と同じだから書き換えるところも少ないんじゃないか？
2011/09/16(金) 18:41:47.81
「Ivy Bridge」プロセッサ、「OpenCL 」をサポート
2011/09/16(金) 23:43:18.97
OpenCLを使ったおぉっとなるアプリってありますか？
2011/09/17(土) 08:21:50.52
うわぁってなるあぷりならある
2011/09/17(土) 08:36:22.90
たとえば？
2011/09/17(土) 18:08:09.83
まってて、ちょっと調べてくる
2011/09/18(日) 07:57:32.03
つ「http://www.geeks3d.com/20110822/webcl-nokia-extension-for-firefox-6-and-kernel-toy/」
2011/09/23(金) 23:20:48.33
なんでAMDのOpenCLってこんなにサポートが糞なの？
2011/09/24(土) 06:50:26.05
OpenCLに関しちゃAMDが一番マシじゃね。
2011/09/24(土) 15:07:51.08
>>500
ヒトイネ
2011/09/27(火) 22:25:38.34
>500
バグレポートしたら、ちゃんと対応してもらえたよ？
2011/10/03(月) 02:47:52.78
もしかして関数呼び出しはOpenCL Cで記述出来ない．．．？
cudaでいう__device__のように記述出来るかと思ったらそんなことはないのかひどいぞこれは
2011/10/03(月) 02:54:10.32
そんなことはなかったすみませんでした
2011/10/03(月) 22:18:02.72
>>503
バグレポートなんか出来るの？
2011/10/04(火) 00:03:56.22
コンパイル時に死ぬバグだけどコンパイラのサポート窓口は分からなかったので
KernelAnalyzerのAboutにあるメアドに「KernelAnalyzerが死ぬんだけど」
って送った。（嘘はついてないつもり）
コンパイラチームに転送してくれて11.10/2.6で直るって連絡来たよ。
実行時に死ぬのとか結果が変なのもForumに書けば結構みてくれるよ。
2011/10/05(水) 23:41:21.79
>>507
Forumに書いても放置が多いみたいで困る。。
509デフォルトの名無しさん
垢版 |
2011/10/06(木) 05:23:23.41
俺が今書いてるJavaプログラムの中で
２つのbyte配列に対する１００〜５万回くらいの論理積（単にforループで２つの配列の論理積をとる）が頻繁に発生するんだけど
そういうとこでOpenCL使ったら高速化期待できる？
まとめて１００万回くらいならＧＰＵが速そうなんだけど
１００回で済むとかだとオーバーヘッドがでかいのかなと思って躊躇してる
2011/10/06(木) 18:38:50.63
aparapiでも使ってみたら？
2011/10/17(月) 07:38:26.26
C++のラッパーを最近使い始めたけどかなり使いやすいね。
とくにメモリ解放が楽になった。
2011/11/02(水) 21:22:05.07
OpenCLは構造体のメモリオブジェクトを作成できますか？

CUDAだったら
typedef struct {
float *num;
} DATA;

DATA data;

cudaMalloc( &data.num, sizeof(float) * 1024 );

みたいにできるんですけど
2011/11/02(水) 23:38:01.05
>>512
例が意味不明
それじゃGPU上に単なるfloat型の配列を確保して、
そのdeviceポインタをCPUの構造体メンバに代入しているだけ。

OpenCLでもGPU上にfloat型の配列をbufferとして確保して、
それをCPUの構造体メンバに代入することは出来る。
メンバ変数の型はfloat*では無く、cl_memだけどね。

でも、その配列自体をGPU側にコピーして、
間接アクセスしようとするとCUDAとOpenCLでは全く違う。

OpenCLではGPU上のポインタは1つのカーネル呼び出し内でしか
一貫性が保証されないから、ポインタを保存しておいて
次のカーネル呼び出しで使うという事が不可能
(cl_memはハンドルに過ぎず、OpenCLのランタイムは
GPU上のオブジェクトを再配置する可能性があるから)
まあ、やるとしたら大きなbufferをメモリプールとして生成して、offsetをポインタの代わりに保存
後でアクセスするときにはbufferを引数に渡して、それにoffset足した位置を触るという
形にするしかない。

まあ、CUDAでも実際にはこの方がCPUとGPUで同じデータを扱えるし、
性能面でも悪くないやり方なんだけど。
2011/11/03(木) 01:23:08.95
ラデ外付けGPUに大量にデータを送りたいんだけど、1/4までって制限どうにかならないの？256MBまでしか送れん
2011/11/03(木) 10:05:29.53
いや、1/4なのではなく 256MBまでという制限。
sizeof(float4) * 4096 * 4096.
2011/11/03(木) 10:38:23.81
>>515
これって変えられないの？
2011/11/04(金) 11:46:02.78
馬鹿には無理
2011/11/07(月) 02:56:58.68
並列化させるのも結構苦労するよね。
簡単な演算ならいいけど、データ依存がちょっとでも複雑になると
動かすカーネルの順番とか数とか
気にしなきゃいけないことｲﾊﾟｰｲ
2011/11/15(火) 23:57:17.03
OpenCL 1.2
http://www.khronos.org/news/press/releases/khronos-releases-opencl-1.2-specification
2011/11/16(水) 00:44:39.04
NVIDIAのドライバがカオスになるな
ｇｄｇｄの果てに漸く1.0対応が落ち着いたと思ったら
28x世代の1.1対応でまたおかしくなってまだ終息してないのに
2011/11/16(水) 01:54:02.46
彼らにはCUDAがあるからなぁ。。
頑張る必要が無いのだろう。
2011/11/19(土) 02:59:11.80
nVidiaはどこに向かっているんだ…

OpenACC : 新しい並列コンピューティングのためのプログラミング環境
http://www.shader.jp/?p=466
2011/11/19(土) 19:26:18.05
それはOpenMPのGPU版みたいなものなので、CUDAやOpenCLとは衝突しない
2011/11/19(土) 20:50:05.10
AlteraがFPGAでOpenCLを、とか言い出してて面白そうな感じ。
2011/11/19(土) 21:07:31.87
C++AMPのNVIDIA版ぽいね
まぁC++AMPはNVIDIAでも動くけど

インライン記述の世代でどれが主導権を握れるかは気になるところ
2011/11/19(土) 23:38:32.61
MSは最終的にC++AMPをオープンにすると言ってはいるが、結局オプソ系コミュニティがどう動くかだな
CPUの並列ライブラリはMSはOpenMPからPPL推し、オプソ系は主にTBBと別れてしまっているので、
GPUではどうにかして歩調を揃えてもらいたいところ

ただ、一応オープン化を標榜するC++ AMPに介入するわけでもなくかといってCUDAの様に自社GPU専用に囲い込むわけでもなく、
立ち位置の被るオープン規格を立ち上げたNVIDIAの意図が分からんといえば分からん
GPUに全てを賭けるメーカーとしては握れる手綱は全て握っておきたい、という事なのかな
2011/11/20(日) 01:50:41.74
>>523
使う側は思いっきり衝突するだろ
2011/11/20(日) 02:32:04.15
>>522
PGI Accelerator が元になっているのかな?
http://www.softek.co.jp/SPG/Pgi/Accel/index.html
2011/11/22(火) 14:27:24.23
インタビューで簡単にOpenACCについてふれてる
http://insidehpc.com/2011/11/21/cuda-reaches-5th-birthday-openacc-ramps-up/
2011/11/22(火) 23:50:42.88
カーネルの実行順位はイベントで指定できる
OpenCLのバイナリコンパイルと読み込みうまくできない・・・

バイナリなしだと環境自由になるけど
ソース丸出しになるから計算高速化くらいしか使い道ないね
2011/11/25(金) 00:37:00.99
1.2の新機能
http://www.streamcomputing.eu/blog/2011-11-19/difference-between-opencl-1-2-and-1-1/
2011/11/25(金) 22:22:16.51
>>531
ダイナミックに追加が変更があったのって、DirectX関係だけだな。。。
2011/11/28(月) 10:27:25.45
なんかもーGPUメーカーは独自に動いてるし、開発する気無いだろｗ
2011/11/29(火) 19:35:56.69
intel CPU制限多すぎ・・・SSE4.1対応って書いておいてくれよ・・・
core2 quad全部対応してるかのようなのはやめてほしい・・・
XPで使えないのはちょっと困る・・・
2011/11/29(火) 23:11:26.28
SSE見逃してたごめん
2011/11/30(水) 00:12:08.31
SSEぐらいOpenCL使わなくってもいいじゃん
2011/11/30(水) 00:50:42.75
>>536
OpenCLの機能実現するのに都合がいい命令がSSE4.1にあるからインテルの開発ツールはSSE4.1対応の世代以降でないと使えないんだよ
どうせ普及する頃には古い世代のCPUいなくなってるよねって方針なんだろ
2011/11/30(水) 00:58:02.95
GPUが使えない環境での互換性用と割り切って広くサポートしてくれたほうがまだ有用なのにな
2011/12/05(月) 06:55:31.93
構造体そのままカーネルに放り投げられない・・・
x.yに配列分けなくちゃだめか・・・
GPU正直CUDAのほうが楽だよね・・・
CPUとGPU同時並列に魅力感じてたけど
両方ともスレッドが少なすぎて・・・
2011/12/05(月) 09:48:13.49
なにいってるんだ？
2011/12/05(月) 11:51:31.75
構造体 a
int x;
int y;
の配列をそのままメモリバッファにコピーしても
うまくいかなかったので・・・
AMDのカーネルアナライザーでエラーがでてたんです
if(a[id].x-a[id].y){}
結局配列分割してコピーしました・・・
2011/12/05(月) 12:04:37.78
組み込みのint2でいいじゃない。
2011/12/05(月) 16:36:44.62
そうですね
それで組み込んでみます
544542
垢版 |
2011/12/05(月) 21:49:58.78
>>543
まず無いだろうけど、別スレッドで
同じ添え字のxとyをバラバラに更新すると嵌るよ。
545543
垢版 |
2011/12/06(火) 04:05:54.65
>>544
xを固定してyの値すべて計算して次のxへという
九九を生成するようなマニアックな使い方をしてるんです・・・
546542
垢版 |
2011/12/06(火) 07:51:19.62
それなら分けた方がいいと思う。
547543
垢版 |
2011/12/06(火) 09:23:20.52
そうしてみます！
2011/12/07(水) 22:41:57.13
ふう・・・ついに完成しました
ちまちま25%使うよりフルロードはいいですね
ただ、オーバーヘッドがいくらあるからはわかりませんが・・・
2011/12/08(木) 07:31:00.39
256MBの制限に引っかかって処理が止まるorz
これの上限増やせないのか?
2011/12/08(木) 10:09:58.25
>>549
グローバルメモリが足りない・・・のかな？
ttp://www.ozone3d.net/gpu_caps_viewer/
これで確認してみて

それかワークスレッドの設定が悪いのかも
2011/12/08(木) 11:03:39.84
>>550
なんでダメか分かった、ありがとう
ローカルメモリ32kbしかないのに2mb使おうとしてたwww
2011/12/08(木) 14:28:54.91
>>551
使い捨てる変数の宣言くらいがちょうどいいよ
NVIDIAはローカルメモリ使わないと倍ぐらい遅くなるけど
他はそうでもないから(OpenCL入門に比較があったよ)
グローバルで問題ないと思うよ
2011/12/15(木) 20:39:38.90
http://www.4gamer.net/games/076/G007660/20111214033/

おまいらこれからもOpenCLにしがみついて行くの？
2011/12/15(木) 20:44:38.54
CUDAはGPUしかできないけど一度により多くの処理ができる
OpenCLはCPUとGPUを同時並列処理ができるのが魅力
どっちにも特化した特性があるからプログラムしだいだよ
2011/12/15(木) 22:16:31.17
GPUには留まらない図が載ってるよ、ヌフォの所に
2011/12/15(木) 22:40:09.72
もうGPU固有じゃなくなってきたのね・・・
カーネル丸出しのCLは論文すら少ない・・・
2011/12/26(月) 03:28:48.61
引数の渡し方が面倒なんだよなー
思わぬところでバグが出たりする
2012/01/18(水) 10:01:49.46
openclを実用的に使うにはどんな環境がおすすめでしょうか
2012/01/18(水) 10:11:34.51
いつでもCUDAに逃げられる環境
2012/01/18(水) 22:05:36.88
nVidiaってこと？
2012/01/19(木) 09:51:18.68
グラフィックボードは倍精度不動小数点数（cl_amd_fp64）が利用可能な
Radeon HD 7900、6900、5900/5800 シリーズがおすすめ

http://blog.tommy6.net/archives/74
2012/01/19(木) 16:52:31.76
宗教戦争が始まりそうだな
2012/01/19(木) 17:05:38.23
>>561
紹介thx
2012/01/19(木) 18:25:09.26
どうせおまえらめんどくさい組み込み関数なんか使わないだろってことか？
2012/01/23(月) 19:34:34.31
GCNでアセンブリが変わったのか、動かなくなってしまった・・・
2012/01/25(水) 02:58:13.00
変わる以前にGCN対応ドライバ＆ランタイムまだ出てないでしょ
CCC12.1Previewが出回ってるからそれ付属のSDKランタイムなら動くんじゃないの
2012/01/25(水) 03:08:02.37
Kernel AnalyzerもTahiti対応版出てないしな
568デフォルトの名無しさん
垢版 |
2012/02/01(水) 22:06:59.00
RADEONとGeforce、ガチンコ対決ではごっちが速いの？
2012/02/01(水) 22:33:40.43
ごっちかな
2012/02/01(水) 23:52:53.73
そもそもガチンコな応用でOpenCLってまともに実績あるの?
2012/02/02(木) 00:02:24.64
無ければなんだっての？
■ このスレッドは過去ログ倉庫に格納されています
5ちゃんねるの広告が気に入らない場合は、こちらをクリックしてください。
ニューススポーツなんでも実況