- 1 名前：a36 ◆K0BqlCB3.k [2008/12/10(水) 15:38:25 .net]
- さてついにOpenCLの仕様が公開されました。
www.khronos.org/opencl/
公式ページにはAPIのヘッダファイルが公開されており、
まだ実際に動かす事はできないもののプログラミングすることは可能となっています。
ということで、公開に先んじてプログラミングを始めてしまいましょう。
- 488 名前：デフォルトの名無しさん mailto:sage [2011/08/09(火) 01:37:46.04 .net]
- >>428
門外不出のノウハウ＝世間一般には知られていない。
つまり各分野
- 489 名前：の研究者すら出し抜く大発見ってことだぞｗ []
- [ここ壊れてます]
- 490 名前：デフォルトの名無しさん mailto:sage [2011/08/11(木) 20:59:03.48 .net]
- >>484
ノウハウが全て開示されているとでも思っているのか？
- 491 名前：デフォルトの名無しさん mailto:sage [2011/08/11(木) 22:05:51.96 .net]
- 世間一般に知られていなければ十分じゃん
- 492 名前：デフォルトの名無しさん mailto:sage [2011/08/12(金) 21:31:49.60 .net]
- 世間一般に知られないようにするためにカーネルを晒したくないわけだろ
- 493 名前：デフォルトの名無しさん mailto:sage [2011/08/14(日) 01:34:57.52 .net]
- >>473
CAL ILコンパイラ自体はいまだに新命令追加されているよ。
64bitINT除算とかはドキュメント化までされている。
ドキュメントにはなっていないけど、wave_idとかは
Appendix Bにはこっそり追加されて、実際テキストで
1OP命令として書けば使える。
でもBFI_INTは置いてけぼり
- 494 名前：デフォルトの名無しさん mailto:sage [2011/08/14(日) 09:57:57.76 .net]
- マスクが定数なら、iand,iand,ior で BFI_INT になるとおもうけど、
そうじゃないときになー。
- 495 名前：デフォルトの名無しさん mailto:sage [2011/08/31(水) 17:10:54.28 .net]
- 初心者です。
Cで書かれた演算プログラムを高速化させたいと思っています。
この元のプログラムをホストプログラムとして、OpenCL用に書き換えて、高速化させたい演算部分を引っこ抜いてカーネルプログラムとして書き換えるというやり方でいいのでしょうか？
また、このようなCのプログラムをOpenCL用に書き換える作業を解説しているサイトや書籍はあるのでしょうか？
当方LINUXでのプログラミングをしております。
- 496 名前：デフォルトの名無しさん mailto:sage [2011/08/31(水) 21:06:58.51 .net]
- 並列化はアルゴリズムをかなり選ぶよ。GPGPUなんて尚更。
まず普通にマルチスレッドで並列化してみたら?
- 497 名前：デフォルトの名無しさん mailto:sage [2011/09/15(木) 10:02:24.21 .net]
- opencl用の自動並列化コンパイラってある？
- 498 名前：デフォルトの名無しさん mailto:sage [2011/09/16(金) 05:07:07.00 .net]
- >>490
そのやりかたでいいさ。
2重ループを1重ループ2つにわけるとか。
あとはほとんどC言語と同じだから書き換えるところも少ないんじゃないか？
- 499 名前：デフォルトの名無しさん mailto:sage [2011/09/16(金) 18:41:47.81 .net]
- 「Ivy Bridge」プロセッサ、「OpenCL 」をサポート
- 500 名前：デフォルトの名無しさん mailto:sage [2011/09/16(金) 23:43:18.97 .net]
- OpenCLを使ったおぉっとなるアプリってありますか？
- 501 名前：デフォルトの名無しさん mailto:sage [2011/09/17(土) 08:21:50.52 .net]
- うわぁってなるあぷりならある
- 502 名前：デフォルトの名無しさん mailto:sage [2011/09/17(土) 08:36:22.90 .net]
- たとえば？
- 503 名前：デフォルトの名無しさん mailto:sage [2011/09/17(土) 18:08:09.83 .net]
- まってて、ちょっと調べてくる
- 504 名前：デフォルトの名無しさん mailto:sage [2011/09/18(日) 07:57:32.03 .net]
- つ「www.geeks3d.com/20110822/webcl-nokia-extension-for-firefox-6-and-kernel-toy/」
- 505 名前：デフォルトの名無しさん mailto:sage [2011/09/23(金) 23:20:48.33 .net]
- なんでAMDのOpenCLってこんなにサポートが糞なの？
- 506 名前：デフォルトの名無しさん mailto:sage [2011/09/24(土) 06:50:26.05 .net]
- OpenCLに関しちゃAMDが一番マシじゃね。
- 507 名前：デフォルトの名無しさん mailto:sage [2011/09/24(土) 15:07:51.08 .net]
- >>500
ヒトイネ
- 508 名前：デフォルトの名無しさん mailto:sage [2011/09/27(火) 22:25:38.34 .net]
- >500
バグレポートしたら、ちゃんと対応してもらえたよ？
- 509 名前：デフォルトの名無しさん mailto:sage [2011/10/03(月) 02:47:52.78 .net]
- もしかして関数呼び出しはOpenCL Cで記述出来ない．．．？
cudaでいう__device__のように記述出来るかと思ったらそんなことはないのかひどいぞこれは
- 510 名前：デフォルトの名無しさん mailto:sage [2011/10/03(月) 02:54:10.32 .net]
- そんなことはなかったすみませんでした
- 511 名前：デフォルトの名無しさん mailto:sage [2011/10/03(月) 22:18:02.72 .net]
- >>503
バグレポートなんか出来るの？
- 512 名前：デフォルトの名無しさん mailto:sage [2011/10/04(火) 00:03:56.22 .net]
- コンパイル時に死ぬバグだけどコンパイラのサポート窓口は分からなかったので
KernelAnalyzerのAboutにあるメアドに「KernelAnalyzerが死ぬんだけど」
って送った。（嘘はついてないつもり）
コンパイラチームに転送してくれて11.10/2.6で直るって連絡来たよ。
実行時に死ぬのとか結果が変なのもForumに書けば結構みてくれるよ。
- 513 名前：デフォルトの名無しさん mailto:sage [2011/10/05(水) 23:41:21.79 .net]
- >>507
Forumに書いても放置が多いみたいで困る。。
- 514 名前：デフォルトの名無しさん [2011/10/06(木) 05:23:23.41 .net]
- 俺が今書いてるJavaプログラムの中で
２つのbyte配列に対する１００〜５万回くらいの論理積（単にforループで２つの配列の論理積をとる）が頻繁に発生するんだけど
そういうとこでOpenCL使ったら高速化期待できる？
まとめて１００万回くらいならＧＰＵが速そうなんだけど
１００回で済むとかだとオーバーヘッドがでかいのかなと思って躊躇してる
- 515 名前：デフォルトの名無しさん mailto:sage [2011/10/06(木) 18:38:50.63 .net]
- aparapiでも使ってみたら？
- 516 名前：デフォルトの名無しさん mailto:sage [2011/10/17(月) 07:38:26.26 .net]
- C++のラッパーを最近使い始めたけどかなり使いやすいね。
とくにメモリ解放が楽になった。
- 517 名前：デフォルトの名無しさん mailto:sage [2011/11/02(水) 21:22:05.07 .net]
- OpenCLは構造体のメモリオブジェクトを作成できますか？
CUDAだったら
typedef struct {
float *num;
} DATA;
DATA data;
cudaMalloc( &data.num, sizeof(float) * 1024 );
みたいにできるんですけど
- 518 名前：デフォルトの名無しさん mailto:sage [2011/11/02(水) 23:38:01.05 .net]
- >>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で同じデータを扱えるし、
性能面でも悪くないやり方なんだけど。
- 519 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 01:23:08.95 .net]
- ラデ外付けGPUに大量にデータを送りたいんだけど、1/4までって制限どうにかならないの？256MBまでしか送れん
- 520 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 10:05:29.53 .net]
- いや、1/4なのではなく 256MBまでという制限。
sizeof(float4) * 4096 * 4096.
- 521 名前：デフォルトの名無しさん mailto:sage [2011/11/03(木) 10:38:23.81 .net]
- >>515
これって変えられないの？
- 522 名前：デフォルトの名無しさん mailto:sage [2011/11/04(金) 11:46:02.78 .net]
- 馬鹿には無理
- 523 名前：デフォルトの名無しさん mailto:sage [2011/11/07(月) 02:56:58.68 .net]
- 並列化させるのも結構苦労するよね。
簡単な演算ならいいけど、データ依存がちょっとでも複雑になると
動かすカーネルの順番とか数とか
気にしなきゃいけないことｲﾊﾟｰｲ
- 524 名前：デフォルトの名無しさん mailto:sage [2011/11/15(火) 23:57:17.03 .net]
- OpenCL 1.2
www.khronos.org/news/press/releases/khronos-releases-opencl-1.2-specification
- 525 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 00:44:39.04 .net]
- NVIDIAのドライバがカオスになるな
ｇｄｇｄの果てに漸く1.0対応が落ち着いたと思ったら
28x世代の1.1対応でまたおかしくなってまだ終息してないのに
- 526 名前：デフォルトの名無しさん mailto:sage [2011/11/16(水) 01:54:02.46 .net]
- 彼らにはCUDAがあるからなぁ。。
頑張る必要が無いのだろう。
- 527 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 02:59:11.80 .net]
- nVidiaはどこに向かっているんだ…
OpenACC : 新しい並列コンピューティングのためのプログラミング環境
www.shader.jp/?p=466
- 528 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 19:26:18.05 .net]
- それはOpenMPのGPU版みたいなものなので、CUDAやOpenCLとは衝突しない
- 529 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 20:50:05.10 .net]
- AlteraがFPGAでOpenCLを、とか言い出してて面白そうな感じ。
- 530 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 21:07:31.87 .net]
- C++AMPのNVIDIA版ぽいね
まぁC++AMPはNVIDIAでも動くけど
インライン記述の世代でどれが主導権を握れるかは気になるところ
- 531 名前：デフォルトの名無しさん mailto:sage [2011/11/19(土) 23:38:32.61 .net]
- MSは最終的にC++AMPをオープンにすると言ってはいるが、結局オプソ系コミュニティがどう動くかだな
CPUの並列ライブラリはMSはOpenMPからPPL推し、オプソ系は主にTBBと別れてしまっているので、
GPUではどうにかして歩調を揃えてもらいたいところ
ただ、一応オープン化を標榜するC++ AMPに介入するわけでもなくかといってCUDAの様に自社GPU専用に囲い込むわけでもなく、
立ち位置の被るオープン規格を立ち上げたNVIDIAの意図が分からんといえば分からん
GPUに全てを賭けるメーカーとしては握れる手綱は全て握っておきたい、という事なのかな
- 532 名前：デフォルトの名無しさん mailto:sage [2011/11/20(日) 01:50:41.74 .net]
- >>523
使う側は思いっきり衝突するだろ
- 533 名前：デフォルトの名無しさん mailto:sage [2011/11/20(日) 02:32:04.15 .net]
- >>522
PGI Accelerator が元になっているのかな?
www.softek.co.jp/SPG/Pgi/Accel/index.html
- 534 名前：デフォルトの名無しさん mailto:sage [2011/11/22(火) 14:27:24.23 .net]
- インタビューで簡単にOpenACCについてふれてる
insidehpc.com/2011/11/21/cuda-reaches-5th-birthday-openacc-ramps-up/
- 535 名前：デフォルトの名無しさん mailto:sage [2011/11/22(火) 23:50:42.88 .net]
- カーネルの実行順位はイベントで指定できる
OpenCLのバイナリコンパイルと読み込みうまくできない・・・
バイナリなしだと環境自由になるけど
ソース丸出しになるから計算高速化くらいしか使い道ないね
- 536 名前：デフォルトの名無しさん mailto:sage [2011/11/25(金) 00:37:00.99 .net]
- 1.2の新機能
www.streamcomputing.eu/blog/2011-11-19/difference-between-opencl-1-2-and-1-1/
- 537 名前：デフォルトの名無しさん mailto:sage [2011/11/25(金) 22:22:16.51 .net]
- >>531
ダイナミックに追加が変更があったのって、DirectX関係だけだな。。。
- 538 名前：デフォルトの名無しさん mailto:sage [2011/11/28(月) 10:27:25.45 .net]
- なんかもーGPUメーカーは独自に動いてるし、開発する気無いだろｗ
- 539 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 19:35:56.69 .net]
- intel CPU制限多すぎ・・・SSE4.1対応って書いておいてくれよ・・・
core2 quad全部対応してるかのようなのはやめてほしい・・・
XPで使えないのはちょっと困る・・・
- 540 名前：デフォルトの名無しさん mailto:sage [2011/11/29(火) 23:11:26.28 .net]
- SSE見逃してたごめん
- 541 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 00:12:08.31 .net]
- SSEぐらいOpenCL使わなくってもいいじゃん
- 542 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 00:50:42.75 .net]
- >>536
OpenCLの機能実現するのに都合がいい命令がSSE4.1にあるからインテルの開発ツールはSSE4.1対応の世代以降でないと使えないんだよ
どうせ普及する頃には古い世代のCPUいなくなってるよねって方針なんだろ
- 543 名前：デフォルトの名無しさん mailto:sage [2011/11/30(水) 00:58:02.95 .net]
- GPUが使えない環境での互換性用と割り切って広くサポートしてくれたほうがまだ有用なのにな
- 544 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 06:55:31.93 .net]
- 構造体そのままカーネルに放り投げられない・・・
x.yに配列分けなくちゃだめか・・・
GPU正直CUDAのほうが楽だよね・・・
CPUとGPU同時並列に魅力感じてたけど
両方ともスレッドが少なすぎて・・・
- 545 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 09:48:13.49 .net]
- なにいってるんだ？
- 546 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 11:51:31.75 .net]
- 構造体 a
int x;
int y;
の配列をそのままメモリバッファにコピーしても
うまくいかなかったので・・・
AMDのカーネルアナライザーでエラーがでてたんです
if(a[id].x-a[id].y){}
結局配列分割してコピーしました・・・
- 547 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 12:04:37.78 .net]
- 組み込みのint2でいいじゃない。
- 548 名前：デフォルトの名無しさん mailto:sage [2011/12/05(月) 16:36:44.62 .net]
- そうですね
それで組み込んでみます
- 549 名前：542 mailto:sage [2011/12/05(月) 21:49:58.78 .net]
- >>543
まず無いだろうけど、別スレッドで
同じ添え字のxとyをバラバラに更新すると嵌るよ。
- 550 名前：543 mailto:sage [2011/12/06(火) 04:05:54.65 .net]
- >>544
xを固定してyの値すべて計算して次のxへという
九九を生成するようなマニアックな使い方をしてるんです・・・
- 551 名前：542 mailto:sage [2011/12/06(火) 07:51:19.62 .net]
- それなら分けた方がいいと思う。
- 552 名前：543 mailto:sage [2011/12/06(火) 09:23:20.52 .net]
- そうしてみます！
- 553 名前：デフォルトの名無しさん mailto:sage [2011/12/07(水) 22:41:57.13 .net]
- ふう・・・ついに完成しました
ちまちま25%使うよりフルロードはいいですね
ただ、オーバーヘッドがいくらあるからはわかりませんが・・・
- 554 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 07:31:00.39 .net]
- 256MBの制限に引っかかって処理が止まるorz
これの上限増やせないのか?
- 555 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 10:09:58.25 .net]
- >>549
グローバルメモリが足りない・・・のかな？
ttp://www.ozone3d.net/gpu_caps_viewer/
これで確認してみて
それかワークスレッドの設定が悪いのかも
- 556 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 11:03:39.84 .net]
- >>550
なんでダメか分かった、ありがとう
ローカルメモリ32kbしかないのに2mb使おうとしてたwww
- 557 名前：デフォルトの名無しさん mailto:sage [2011/12/08(木) 14:28:54.91 .net]
- >>551
使い捨てる変数の宣言くらいがちょうどいいよ
NVIDIAはローカルメモリ使わないと倍ぐらい遅くなるけど
他はそうでもないから(OpenCL入門に比較があったよ)
グローバルで問題ないと思うよ
- 558 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 20:39:38.90 .net]
- www.4gamer.net/games/076/G007660/20111214033/
おまいらこれからもOpenCLにしがみついて行くの？
- 559 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 20:44:38.54 .net]
- CUDAはGPUしかできないけど一度により多くの処理ができる
OpenCLはCPUとGPUを同時並列処理ができるのが魅力
どっちにも特化した特性があるからプログラムしだいだよ
- 560 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 22:16:31.17 .net]
- GPUには留まらない図が載ってるよ、ヌフォの所に
- 561 名前：デフォルトの名無しさん mailto:sage [2011/12/15(木) 22:40:09.72 .net]
- もうGPU固有じゃなくなってきたのね・・・
カーネル丸出しのCLは論文すら少ない・・・
- 562 名前：デフォルトの名無しさん mailto:sage [2011/12/26(月) 03:28:48.61 .net]
- 引数の渡し方が面倒なんだよなー
思わぬところでバグが出たりする
- 563 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 10:01:49.46 .net]
- openclを実用的に使うにはどんな環境がおすすめでしょうか
- 564 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 10:11:34.51 .net]
- いつでもCUDAに逃げられる環境
- 565 名前：デフォルトの名無しさん mailto:sage [2012/01/18(水) 22:05:36.88 .net]
- nVidiaってこと？
- 566 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 09:51:18.68 .net]
- グラフィックボードは倍精度不動小数点数（cl_amd_fp64）が利用可能な
Radeon HD 7900、6900、5900/5800 シリーズがおすすめ
blog.tommy6.net/archives/74
- 567 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 16:52:31.76 .net]
- 宗教戦争が始まりそうだな
- 568 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 17:05:38.23 .net]
- >>561
紹介thx
- 569 名前：デフォルトの名無しさん mailto:sage [2012/01/19(木) 18:25:09.26 .net]
- どうせおまえらめんどくさい組み込み関数なんか使わないだろってことか？
- 570 名前：デフォルトの名無しさん mailto:sage [2012/01/23(月) 19:34:34.31 .net]
- GCNでアセンブリが変わったのか、動かなくなってしまった・・・
- 571 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 02:58:13.00 .net]
- 変わる以前にGCN対応ドライバ＆ランタイムまだ出てないでしょ
CCC12.1Previewが出回ってるからそれ付属のSDKランタイムなら動くんじゃないの
- 572 名前：デフォルトの名無しさん mailto:sage [2012/01/25(水) 03:08:02.37 .net]
- Kernel AnalyzerもTahiti対応版出てないしな
- 573 名前：デフォルトの名無しさん [2012/02/01(水) 22:06:59.00 .net]
- RADEONとGeforce、ガチンコ対決ではごっちが速いの？
- 574 名前：デフォルトの名無しさん mailto:sage [2012/02/01(水) 22:33:40.43 .net]
- ごっちかな
- 575 名前：デフォルトの名無しさん mailto:sage [2012/02/01(水) 23:52:53.73 .net]
- そもそもガチンコな応用でOpenCLってまともに実績あるの?
- 576 名前：デフォルトの名無しさん mailto:sage [2012/02/02(木) 00:02:24.64 .net]
- 無ければなんだっての？
- 577 名前：デフォルトの名無しさん mailto:sage [2012/02/02(木) 15:35:32.54 .net]
- いや別に問題ない
- 578 名前：デフォルトの名無しさん [2012/02/10(金) 23:47:32.03 .net]
- とりあえず齧る分には公式のプログラミングガイド買っとけばおｋ？
- 579 名前：デフォルトの名無しさん mailto:sage [2012/02/13(月) 02:50:23.08 .net]
- OK
C言語と同じ
- 580 名前：デフォルトの名無しさん mailto:sage [2012/02/22(水) 23:40:42.01 .net]
- サブルーチンみたいにカーネルに直接引数渡して処理できればなぁ
アドレスを渡す時とかすげえ面倒。
- 581 名前：デフォルトの名無しさん mailto:sage [2012/02/23(木) 01:38:17.77 .net]
- そんな気軽にあちこちで使うようなもんじゃないっしょ
- 582 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 16:45:33.65 .net]
- カーネルソース丸出しか
特定デバイス用にコンパイルしておかないと
いけないのがなあ
- 583 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 16:57:57.20 .net]
- 大したことないやつほど隠したがる銭湯の法則
- 584 名前：デフォルトの名無しさん mailto:sage [2012/03/08(木) 22:38:56.05 .net]
- BOINC的なもので使うと不正対策必要で、既知の答えも一緒に計算させたりして対策するんだけどエコじゃないて悩ましいよ
- 585 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 01:23:31.48 .net]
- それなら同じWU他のやつに配って結果一致したやつだけ採用、
合わなかったら更に配布して一致した方採用でいいんでないかね
どのみち普通にやったって計算エラーになるのもいるし
- 586 名前：デフォルトの名無しさん mailto:sage [2012/03/11(日) 02:45:56.59 .net]
- つーかカーネルがソースコードだから
改竄される恐れがなんて懸念するような
プログラムなら、どんな形態で配るにせよ
ディスアセンブルされりゃ同じだって。
- 587 名前：デフォルトの名無しさん mailto:sage [2012/03/17(土) 16:17:03.96 .net]
- 個人PCでスパコンの1/1000の計算速度（10テラFLOPS）だってよ。スゲーな。
dualsocketworld.blog134.fc2.com/blog-date-201202.html
- 588 名前：デフォルトの名無しさん mailto:sage [2012/03/18(日) 17:23:23.96 .net]
- 10TFLOPSっていうと10年前のスパコンの性能と同じくらいだな
|
 |
