OpenCLプログラミング#1

**a36** ◆K0BqlCB3.k · 2008/12/10(水) 15:38:25

さてついにOpenCLの仕様が公開されました。

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

公式ページにはAPIのヘッダファイルが公開されており、
まだ実際に動かす事はできないもののプログラミングすることは可能となっています。
ということで、公開に先んじてプログラミングを始めてしまいましょう。

**デフォルトの名無しさん** · 2010/05/20(木) 22:39:47

個人的な覚書

ATI Stream 2.1のcl_platform.hはこのままだとmingw-gccで使えない。
38行目の#include <stdint.h>を削除して
77行目 #else /* !_WIN32*/の次の行に入れる必要あり。

khrnosの配っているcl.hpp　(\date $Date: 2010-04-23 10:16:50 -0500 (Fri, 23 Apr 2010) $は
namespace cl　のなかでstd::pairと呼び出してる箇所を全て ::std::pairと修正、
#if defined(__CL_ENABVLE_EXCEPTIONS)
#include <exception>
#endif
をnamespace clの外側に移動する必要あり。

**デフォルトの名無しさん** · 2010/06/03(木) 07:06:02

ワープって何よ。
亜空間跳躍となにか関係あるのか?

**デフォルトの名無しさん** · 2010/06/03(木) 23:01:30

OpenCLはワープサイズやウェイブフロントサイズをとってこれないのが難点だな。
これが取れないとポータビリティのある最適化はできないよ～。

**デフォルトの名無しさん** · 2010/06/04(金) 13:26:57

ワープとかの概念が無用なプラットフォームも OpenCL の対象範囲。
どうせハードウェアごとの最適化が入る部分は、CL がどうこうする部分じゃないよ。
その辺は #ifdef で。

**デフォルトの名無しさん** · 2010/06/04(金) 22:55:03

ワープサイズに依存しない機種の場合はワープサイズ1でも返してくれればいいんだけど。

**デフォルトの名無しさん** · 2010/07/01(木) 23:02:44

コードのコンパイル、デバイスの生成、データ転送に喰われてスケールしないんだが。
ありがた味のわかるサンプル教えてくれ。

**デフォルトの名無しさん** · 2010/07/01(木) 23:05:51

コンパイルやデバイスの生成やデータ転送と比べものにならないくらい長時間のデータを計算させる。

**デフォルトの名無しさん** · 2010/07/03(土) 00:48:14

openclでGPGPUするかどうかの判定ラインって何？
やってみないとわからんのか。

**デフォルトの名無しさん** · 2010/07/03(土) 12:42:13

計算量がNの3乗くらいなら有効だと思うよ。
ループが3重にネストしているような奴。

**デフォルトの名無しさん** · 2010/07/24(土) 22:38:24

水を差すようだけど，OpenCLだからと言ってO(N^3)がO(N^2)に落ちる訳ではないんだから
それだけでは何の基準にもならないよね．
滅茶重い計算だとNも単純に大きくてデバイスローカルなメモリに収まらなかったりするし
対象デバイスのwarp sizeを想定しながらブロック化とか，結局ifdef祭り．

**デフォルトの名無しさん** · 2010/07/24(土) 23:00:12

うるせえだまってろ

**デフォルトの名無しさん** · 2010/07/25(日) 08:00:33

でっていう

**デフォルトの名無しさん** · 2010/09/02(木) 01:36:07

もう少しなんかあっても良いよね。

**デフォルトの名無しさん** · 2010/10/10(日) 01:30:39

文字列処理のサンプルってないですか？

**デフォルトの名無しさん** · 2010/10/24(日) 00:50:26

新しいMacBookPro買ったんだけど
clCreateContextFromTypeでGPU指定したら、そこからCL_SUCCESSが返ってこない・・・
Appleのサイト見てもGeForce GT 330Mはサポートされてるのに

**デフォルトの名無しさん** · 2010/10/24(日) 01:04:56

>>353
nVidiaが公開してるCUDAドライバ入れた？

**デフォルトの名無しさん** · 2010/10/24(日) 01:33:46

>>354
SnowLeopardはOSでサポートされてるんじゃないの？
ちょっとぐぐってきます・・・

**デフォルトの名無しさん** · 2010/10/24(日) 10:40:54

ドライバもなしにハードを動かそうとな？

**デフォルトの名無しさん** · 2010/10/24(日) 16:18:23

http://yanoshin.jp/entry_154.html
いつのバージョンかわからんけど
＞＊必ずToolkit、Driver、SDKの順にインストールすること！（なぜか？）

ドライバだけで、他は必要ないのかな？
http://journal.mycom.co.jp/column/osx/335/index.html
http://developer.apple.com/library/mac/#documentation/Performance/Conceptual/OpenCL_MacProgGuide/OpenCLontheMacPlatform/OpenCLontheMacPlatform.html

**デフォルトの名無しさん** · 2010/10/24(日) 16:21:17

今までRadeonユーザーにOpenCL使ってもらうには
SDKインストールしてもらわないといけなかったんだが、
最新ドライバから、OpenCLランタイムをインストールできるようになったらしい。
めでたいことだ。

**デフォルトの名無しさん** · 2010/10/24(日) 17:39:16

Radeonって全然高速じゃないし使う意味ないよなぁ

**デフォルトの名無しさん** · 2010/10/24(日) 18:35:46

Radeonユーザーの場合はDirectCompute使った方が良いね

**デフォルトの名無しさん** · 2010/10/24(日) 18:41:34

Core i　……　RADEON HD 6xxx
Phenom II　……　GeForce GTX 4xx

**デフォルトの名無しさん** · 2010/10/24(日) 19:48:56

RadeonだとN体 24Gぐらいしかでないしな
Geforceだと500Gぐらい平気で出るんだけどねぇ

**デフォルトの名無しさん** · 2010/10/24(日) 23:22:24

>>356
秀和システムのOpenCL入門を読んだら
Mac：Xcodeあればいい
Win：ドライバダウンロード必要
としか書いてなかったからてっきり。
>>357
インストール順番さっそくミスった
実はMacは初めてだからアンインストールの仕方ぐぐらんと

**363** · 2010/10/27(水) 09:01:28

やっぱり動かない・・・
よく見たらApple英語サイトはGeForceGT３３０M書いてあるけど
日本語のほうは載ってないし
http://www.apple.com/macosx/specs.html
http://www.apple.com/jp/macosx/specs.html

**デフォルトの名無しさん** · 2010/10/27(水) 11:20:00

それ、OSの必要条件じゃん。だから何?

**デフォルトの名無しさん** · 2010/10/29(金) 07:17:55

http://developer.apple.com/library/mac/#samplecode/OpenCL_Hello_World_Example/Introduction/Intro.html
> __kernel square
を
> __kernel void square
に置換。

のサンプルも動かない？

**デフォルトの名無しさん** · 2010/10/29(金) 16:50:26

>>365
ごめん　それは関係なかった
>>366
カーネル以前にclGetDeviceIDsでGPU指定だと止まる。CPUでは動く。。。
試しに>>357の参考してCUDAのサンプルも動かしてみたら
CUDA使用可能なデバイスが見つからないというエラーがでる(´･ω･`)
でも動くのもあって、oceanFFTとかnbodyとかは動いてくれるんだけど

**デフォルトの名無しさん** · 2010/10/29(金) 19:14:57

そりゃ変だな。ちなみにこちらは Radeon HD 5670 の iMac だが普通に動いてる。
CUDAのドライバにSDKは入れてる（入れただけで放置してたが）。

**デフォルトの名無しさん** · 2010/11/06(土) 16:01:55

gccはOpenCLに興味ないみたいだね
代わりにclang使うしかないのか?UNIXでは

**デフォルトの名無しさん** · 2010/11/07(日) 00:02:23

>>369
gccは開発技術レベル落ちまくってて
本業のコンパイラもLLVM連中やMSに
抜き去られてやばいらしい

**デフォルトの名無しさん** · 2010/11/07(日) 00:11:37

つってもGCC死んだらUNIX系が基盤から崩壊するから、
どうにかしなきゃいかんだろうしどうにかするだろ

**デフォルトの名無しさん** · 2010/11/07(日) 00:29:49

マルチアーキ型のgccの最適化処理がx86特化型であるmsのソレ以上だったことなんてあったか？

**デフォルトの名無しさん** · 2010/11/07(日) 00:30:00

GCCってそんな立ち位置だったっけ？
広く愚直にというイメージがあるんだが

**,,・´∀｀・,,）っ-○○○** · 2010/11/07(日) 01:03:36

GNUソフトウェアはいつだって言いだしっぺの法則が適用されるんだぜ。
Cell B.E.だってSCEが自らABIを公開してGCCやLinuxカーネルのパッチをコミットした。

NVIDIAやAMDがGCCにコミットしない限り動かんだろ。

**デフォルトの名無しさん** · 2010/11/07(日) 05:02:21

>>371
もうどうにもならん
ORACLEのせい

**,,・´∀｀・,,）っ-○○○** · 2010/11/07(日) 05:21:45

なんだかんだでHPCなんて所詮はサーバ市場の１％未満のニッチ市場だよ。
そのなかのGPGPU用の言語の一つのために必死になる奴もそうそういない。
なんならRailsやZend使ってWebサービスを構築してるSIerにでもGPGPUの素晴らしさを説いてくれｗ

**デフォルトの名無しさん** · 2010/11/07(日) 10:51:44

>>376
動画のエンコード・デコードはダメなの？？

**デフォルトの名無しさん** · 2010/11/07(日) 13:34:36

ffmpegと同等の画質を達成すればGoogle様が興味を示す

かも

**デフォルトの名無しさん** · 2010/11/07(日) 13:44:30

なんでGPUを使う→低画質という思考になるの？

**デフォルトの名無しさん** · 2010/11/07(日) 14:24:23

馬鹿だから

**,,・´∀｀・,,）っ-○○○** · 2010/11/07(日) 15:08:22

ブロックサイズを細かく分けて並列処理しないとGPUのスループットの強みを生かせないから。
CPUで処理するレベルの高画質エンコードは構造的に不向き。せいぜいフィルタだけやらせておけというレベル。

**デフォルトの名無しさん** · 2010/11/07(日) 20:04:27

ブロックサイズ1ピクセル単位とかにでもしないとGPUの優位なんて出てこないんだろうなｗ

**デフォルトの名無しさん** · 2010/11/17(水) 00:02:22

IntelがCoreプロセッサ向けOpenCL SDKのアルファ版を公開: 新 masafumi's Diary
ttp://masafumi.cocolog-nifty.com/masafumis_diary/2010/11/intelcoreopencl.html

**デフォルトの名無しさん** · 2010/11/17(水) 01:07:04

Windows Onlyか

**デフォルトの名無しさん** · 2010/11/17(水) 15:57:12

Core2Duoなのに入れられない！と思ったら45nm版Core2以降かよ…SSE4使うのな

**デフォルトの名無しさん** · 2010/11/17(水) 19:24:56

そろそろQ6600に引導渡してやるか・・・
Cougar Pointの仕様が色々と半端なので、GPUだけ強化してIvy Bridgeまで引っ張るつもりだったんだが

**デフォルトの名無しさん** · 2010/11/21(日) 00:35:46

来年からAMDが大攻勢、NVは苦しくなるな。
もうグラフィックボードすら不要になってくる。

**デフォルトの名無しさん** · 2010/11/21(日) 00:38:40

CPUで使う場合も、OpenCL的なメモリコピーを挟まないといけないのでしょうか？

**デフォルトの名無しさん** · 2010/11/27(土) 07:53:49

OpenCLって、いつになったら動かせるようになるの？
組み込み向けにGPGPUしたいんだけど、CUDAだと発熱すごすぎグラフィックボード
がどうしても必要なんで、オンボードGPUをOpenCLで動かしたいのだが。

**デフォルトの名無しさん** · 2010/11/27(土) 08:38:35

まるでオンボードGPUだと発熱しないと言わんばかりの意味不明な発言。

**デフォルトの名無しさん** · 2010/11/27(土) 10:03:51

信頼性は違うだろ？

**デフォルトの名無しさん** · 2010/11/27(土) 13:28:35

オンボードのチップの方が信頼性が高いの?
マザボ丸ごと道連れになるくらいならカード交換で済む方がいい気もするんだが。
今までの経験からすると、寧ろ電源が死ぬことのほうが多いけどね。

**デフォルトの名無しさん** · 2010/11/27(土) 14:29:27

RadeonでもGeForceでもPhenomでもCore 2(45nm以下)以上でも動くよ>>OpenCL

**,,・´∀｀・,,）っ-○○○** · 2010/11/27(土) 14:34:16

SSE4.1以上に限定する意味は一応あるのね。
insertps/extractpsがgather/scatterとして、vlendvpsがプレディケートに使えるから
SPMDのカーネルモデルを容易に実現できる。

**デフォルトの名無しさん** · 2010/11/29(月) 14:16:12

OPENCLってグラフィックカードの計算につかうCPUで
普通のプログラミングも計算するってことですか？

**デフォルトの名無しさん** · 2010/11/29(月) 19:38:51

日本語でおｋ

**デフォルトの名無しさん** · 2010/11/29(月) 22:13:20

>>395
コストが無駄になるので特定の用途向けの計算に使います。

**デフォルトの名無しさん** · 2010/12/01(水) 11:07:07

じゃあ、グラフィックカードの付いてないノートパソコンでは無意味ということだな。

**デフォルトの名無しさん** · 2010/12/02(木) 01:56:33

C言語っぽく書いて、コンパイラがうまいことSIMDに変換してくれるから意味あるよ。

**デフォルトの名無しさん** · 2010/12/04(土) 11:20:13

GPUでmalloc作った論文どこだっけ？

**デフォルトの名無しさん** · 2010/12/18(土) 09:11:15

外付けと内蔵グラフィック切り替えるタイプのノートでopenCL使うとどうなるの
NVkernelそこまで対応してるのかな？

**デフォルトの名無しさん** · 2010/12/25(土) 19:51:36

全ワークアイテムでは同期とれないんだね。ワークグループ内のみか。
さて、どうやって同期とるか・・・

**デフォルトの名無しさん** · 2010/12/31(金) 18:12:58

OpenCLを始めようと思ってサンプルコードを動かしてみたんですが
コンテキストを作るところで躓いてしまいました。
何がいけないんでしょうか？

具体的には
cl_int status;
cl_context context = clCreateContextFromType(NULL, CL_DEVICE_TYPE_GPU, NULL, NULL, &status);
とやると、status に CL_INVALID_PLATFORM が返ってきます。

clGetPlatformInfo 関数は、たとえば CL_PLATFORM_VERSION を引数にして実行すると
"OpenCL 1.0 CUDA 3.2.1" という文字列が返ってくるので
OpenCLドライバが入ってないわけではなさそうなんですが・・・

環境はCore2Duo E8400、GeForce9600GT、Windows7 32bitです。

**デフォルトの名無しさん** · 2010/12/31(金) 18:39:37

第一引数がNULLだとそうなるのがある
ならないのがあるのかは知らん
秀和の本読んでると思うけど公式から修正されたソース落としてこれるからそれを参考にしてくれ

今規制中で携帯からコード打つ気にならんｗ

**403** · 2010/12/31(金) 20:39:19

>>404
秀和のサイトの修正済サンプルコードを参考にしたら、うまく動きました。
第一引数にNULLではなくコンテキストのプロパティを与えればいいみたいですね。

秀和の本で勉強してると言い当てられてビックリですｗ
この本わかりやすいですね。

大学の研究で何千枚の画像相関やってるんで、OpenCLで夢がひろがりんぐ。

**デフォルトの名無しさん** · 2011/01/01(土) 10:51:11

グラフィックスドゥライヴァのヴァージョンは？

**404** · 2011/01/01(土) 16:41:42

http://developer.download.nvidia.com/compute/cuda/3_0/toolkit/docs/NVIDIA_OpenCL_ImplementationNotes_3.0.txt

**デフォルトの名無しさん** · 2011/01/04(火) 05:00:19

openCLを使っています。
CL_DEVICE_TYPE_DEFAULTという値について質問があります。
この値はプラットフォームの標準デバイスを使うと書いてありました。
どのようにしたら、プラットフォームの標準デバイスがGPUなのか、CPU
なのかを知ることができるのでしょうか。

**デフォルトの名無しさん** · 2011/01/26(水) 11:02:46

Radeon5850と6970を二本さして、FFTの負荷分散させたいと思ってcontextにデバイス２つのせてキューも2本立てたんだけど、
データ転送も並列実行も同時並列的に実行されないで逐次的に実行されてるみたいなんですよ。
コードは1GPUの場合のものをキューだけ変えて単純に2つ並べただけなんですけれども。

OpenCL StramSDKでマルチGPU?実行されてる方いらっしゃいます？

**デフォルトの名無しさん** · 2011/01/26(水) 12:09:39

AMDからメール来た

AMD Accelerated Parallel Processing (APP) SDK | AMD Developer Central
http://developer.amd.com/gpu/AMDAPPSDK/Pages/default.aspx
OpenCL? H.264/AVC: MainConcept
http://www.mainconcept.com/products/sdks/gpu-acceleration/opencltm-h264avc.html

**409** · 2011/01/28(金) 13:15:14

OpenMP使えばいいんですね．自己解決しましたｗ

**デフォルトの名無しさん** · 2011/01/28(金) 18:00:42

OpenCLはCUDAやTBBのラッパーという理解で宜しいか？

**デフォルトの名無しさん** · 2011/01/31(月) 10:23:07

OpenMPはCPUしか並列化できない
GPUも並列化したいならOpenMPじゃなくてOpenCL使わなきゃだめだろ

**デフォルトの名無しさん** · 2011/01/31(月) 21:48:44

>>413
レスあざす。
もともとOpenCLで組んでて、OpenCLだけでマルチデバイスで動作しないかなぁと思った悩んでいたもんで。一応ホストのスレッドを二本立てたら並列動作しました！

**デフォルトの名無しさん** · 2011/01/31(月) 23:26:33

AMDのH264のSDKってどうなの？

**デフォルトの名無しさん** · 2011/02/01(火) 20:28:57

>>415
んなもんあったっけ？？

**デフォルトの名無しさん** · 2011/02/02(水) 01:05:26

>>416
話解らない奴だな
OpenCL H.264/AVC Encoder SDK Beta

どうなんだって聞いてるんだから答えろやクソが

**デフォルトの名無しさん** · 2011/02/02(水) 15:44:26

>>417
それAMD製じゃないしAMD専用でも無いじゃん

http://www.mainconcept.com/products/sdks/gpu-acceleration/opencltm-h264avc.html

**デフォルトの名無しさん** · 2011/02/02(水) 21:39:51

>>417
ぷｗ

**デフォルトの名無しさん** · 2011/02/03(木) 00:23:39

>>414
OpenCLの関数ってほとんどノンブロッキングだから
OpenMP使わなくても並列に動作するんじゃないの？
試したことないからわからんけど

**デフォルトの名無しさん** · 2011/02/03(木) 00:34:19

>>420
と思ったんですよ。とりあえず同じコンテクストにデバイス２つ乗っけて、キューを二本立ててやってみると計算はされるんだけど、プロファイラーで動作みてみると見事に逐次実行されてて。コンテクスト2つ立ててもダメでした。

ヘテロジニアスなんて謳ってるわりに面倒ですよねぇ^^;

**デフォルトの名無しさん** · 2011/02/04(金) 03:56:08

>>421
OpenCLのコマンドキューはインオーダー実行がデフォだから
CL_QUEUE_OUT_OF_ORDER＿EXEC＿MODE_ENABLEやらないと必ず逐次実行されるよ

**デフォルトの名無しさん** · 2011/02/04(金) 20:08:07

>>422
違うんですねー
OpenCLドライバは何故か初期から、OutOfOrderを付けても逐次実行されるんだわ。。。
いつか対応されるだろうと、一年くらい経過したけどまだ対応されてないとは

**デフォルトの名無しさん** · 2011/02/09(水) 20:10:42

>>421
プロファイラで実行すると逐次実行されるが、プロファイラ上でなければ並列実行される。
ジョブ投入や実行の時間を記録して表示させてみ。

また、同一コンテキストのメモリが二つのデバイスのどちらからでもアクセスできる裏技もあるよ。
最初のアクセスは遅いが、それ以降は早いので、CPU側に読み込んで書き戻すとかしなくてもダミーでメモリをアクセスするジョブを流せばデバイス間のデータ移動ができる。

1年ほど前にデバイス2本同時駆動の実験だけして本チャンの計算プログラムはまだつくってなかったな、そういや。

**デフォルトの名無しさん** · 2011/03/08(火) 07:32:39.91

昨日1日使ってボトルネックをOpenCLに移植したけど、
余計遅くなったわ。やれやれ

**デフォルトの名無しさん** · 2011/03/08(火) 08:28:07.93

ボトルネックを移植するんじゃなくてOpenCLに向いている部分を移植しろよ
または並列性を極大化したりif文などの分岐を全部排除するなどアルゴリズム最適化をしてみれば？

**デフォルトの名無しさん** · 2011/04/07(木) 03:06:36.77

VLIW5のtパイプ埋まらん

**デフォルトの名無しさん** · 2011/04/13(水) 04:08:44.89

深く考えずに2^26スレッドenqueueしたらひどいことになった。

**デフォルトの名無しさん** · 2011/04/21(木) 23:37:32.85

なんでこれ有料なんだ？
ttp://web.me.com/cducommun/CLBuilder/Home.html

**デフォルトの名無しさん** · 2011/04/22(金) 00:28:47.04

Appleが提供しているツールならまだしも第三者が開発しているツールを無料にしろってのは横暴だろ
シェアウェアや市販ソフトを完全否定なお人ですか？

**デフォルトの名無しさん** · 2011/04/23(土) 23:18:19.02

AMD APPのOpenCLについてご存じの方がいたら教えて欲しいんですが、
wavefront=64 って事は、例えばwork_group = 128 の場合 work_itemが64ずつ完全に並列実行って認識で合ってるでしょうか？
（wavefront内のやりとりは同期が必要ないかどうかを知りたいです）

**デフォルトの名無しさん** · 2011/04/24(日) 03:56:08.00

>>431
何の保証もないけど、今のところ大丈夫な気がする。

**デフォルトの名無しさん** · 2011/04/24(日) 05:50:16.61

最大3クロックのタイミング差があると思うんだけど。

**デフォルトの名無しさん** · 2011/05/26(木) 04:58:19.55

XOR_INT*4とMULADD*4が同じ時間食うって納得いかん…

**デフォルトの名無しさん** · 2011/06/14(火) 22:58:32.16

for(d=0;省略){
clSetKernelArg(省略&d);
clEnqueueNDRangekernel(省略);
clEnqueueReadBuffer(省略);
for(省略){読み出したメモリ合算}
}
で1024*dで実行して
読み出して合算して再度実行・・・の形にしてもアドレス違反が起きてしまう・・・
イベントリストもバリア同期も意味なし

GPU実行でアドレス違反は画面だけ真っ黒になってびっくりした
デバックはCPUデバイス使ったほうがいいよ

**デフォルトの名無しさん** · 2011/06/26(日) 01:59:15.81

最近OpenCLに興味もったんですが皆さん楽しいですか？

**デフォルトの名無しさん** · 2011/06/26(日) 14:59:58.32

楽しいとか楽しくないとか問題じゃねえよ

**デフォルトの名無しさん** · 2011/06/26(日) 15:03:37.91

でっかいうんこかどうかだな