前スレ
GPGPU#4
http://hibari.2ch.net/test/read.cgi/tech/1255256230/l50
関連スレ
OpenCLプログラミング#1
http://hibari.2ch.net/test/read.cgi/tech/1228891105/l50
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
http://hibari.2ch.net/test/read.cgi/tech/1271587710/l50
参考リンク
総本山? gpgpu.org
http://www.gpgpu.org/
OpenCL
http://www.khronos.org/opencl/
NVIDIA CUDA
http://developer.nvidia.com/object/cuda.html
ATI Stream
http://developer.amd.com/gpu/ATIStreamSDK/Pages/default.aspx
GPUをCPU的に活用するGPGPUの可能性
http://pcweb.mycom.co.jp/articles/2005/09/06/siggraph2/
＞と言いつつAgeiaの中の人も今じゃAMDにいるからなぁ
＞とんだ詐欺師なのかねあの人
金です。
nvにとっちゃすでに用済みで、要らない子
その分割された小包の集団をどかっとCUDAに押し込んでやると、分割が上手ければ
それなりに速く結果が出る。ただ、GPGPUで効率が出る並列化は簡単ではない。Larrabeeがこけたのもここ。
しかもC++のCUDA方言は不思議挙動だったりで、技術者がCUDAに習熟して十分な速度が
出せるようになるまでの時間を考えると、結構経費がかかる。だから、相当大きな話、というか
CUDAのX86@Intel CPUに対するワットパフォーマンス優位性が技術者の勉強代をカバーできる規模で無いと
わざわざわけわからん方言を勉強したくない。しかも、この方言は、いつまで有効かも怪しい。
だから、ほとんどの用途では、Nehalem-Ex とか、速いCPU乗せたマシンを増やした方が良い。
他のプログラムが、"全部" 速くなりますからねｗ
その対策として出てきたのがCPUのマルチコア化と
グラボのGPGPUとしての活用なわけで・・・
大部分の人には上位CPUなんて必要ないのと同様に
大部分のアプリにもGPGPUなんて必要ない。
６コアもGPGPUも本当に必要な人・アプリが使えばいいだけ
そんなもん無視でいいだろｗ
音楽・動画再生だろうがゲームだろうが大部分のアプリには
高価な上位CPUも高速なGPGPUも必要じゃないだろ。
そこそこヘビーな自分でも4コア(疑似8コア)や
1TFLOPS以上のGPUをフル活用できるのは全PC作業の１割程度だし
やる気無いですよねぇ
条件分岐も少なくないからCPUで計算したほうがいい。
映像編集ソフトですらエフェクト処理がメインでエンコードにはGPGPUが使えなかったりする。
iTunes（映像再生ソフト）にGPGPUとして使うなんて問題外。
大人しくOpenGLやDirect2DなんかでGPUとして活用すべき。
リアルタイムで映像にエフェクト処理を加えながら再生したいなら別だがiTunesの仕事じゃないｗ
PowerDirector？
データ転送がボトルネックなのだから
単位データあたりの演算量を増やさなきゃメリットが無い。
どんだけ複雑になったって、大量に並列実行できればGPGPUにとってアドバンテージがある。
データに対して演算量が少なすぎると転送や処理待ちばかりになってパフォーマンスが上がらない。
だから問題は複雑性よりもデータの相互依存性とデータに対する演算量の少なさ。
もっと端的に言ってしまえば分岐命令の数
この場合は違う。
>14 はそのつもりで発言しているのかも知れないが、>13 は違う。
演算回数は増えるがGPUなら並列数を増やせる感じで
http://www.semiaccurate.com/2010/09/06/what-amds-northern-islands/
ごめんSIって言ってたけど実はNIだったよ。えへ。
だから今度出るのはHD6000ファミリーはNIね。
32nmでNIテープアウトしてたけど40nmで出すよ。
コアは○○な感じで、アンコアは××な感じで強化してるよ。
なんでチップがEvergreenより10-15%大きくなるよ。
リリーススケジュールは10月12日にイベントで25日前後に店頭並ぶよ。
まずはAMDの穴の開いてる$175-250帯のHD6700から始めるよ。
次にHD6800、HD6900、年初にローエンド、28nmまでこのラインナップだよ。
HD6000出たら緑チームはHD5000よりコスト高いのに値下げしなくちゃだし、それでなくても冷め切ったセールスにもろ影響しちゃうよ。
だって、トップエンドは価格維持でHD5000は下がり始めるしね。
Nvidiaの夢と希望を打ち砕いちゃうね。
打つ手もないしね。
AMDはDX11のトーナメント1回戦をHD5000で勝利して、第2回戦もHD6000で勝利しちゃて、Nvidiaには財務的にもパフォーマンスでっかいマージンを取っちゃうよ。
28nmまではNvidiaにチャンスはないね。
窓から投げ捨てるべきでしょうか？
やっぱり実績の無いハウスに委託したのが間違いだったか・・・
弊社ではソースがあればデバッグも承りますw
その辺だけはしっかりしているという・・・
ていうか、ウチ（受け入れ側）のマネージャーが完全に「ドモホルンリンクル」で
どんなゴミを渡されても「努力あるのみ」とかの類の精神論を吐いて話にならないし
どっか、受託開発や納入後の展開方法についての客観的な評価をしてくれる
コンサルタントはないですかね・・・
いつバージョンアップでバイナリが動かなくなってもおかしくないのに
コストの問題だろ？
ソースを要求すると価格が上がる
teslaを用いた計算機を導入しようとしているのですが、i7-980xにするかi7-930にするか
迷っています。
他にもGPGPU用中間コードのコンパイルにも影響するが誤差範囲。
MP4に変換できるソフトってありますか？
あとRADEONのカードってエンコードなら値段と性能みてどれがコスパいいですか？
ソフト板か自作板、DTV板へgo
ここは「ATI技術に対応しててMP4に変換できるソフト」を作る側の板です。
Actor とか Map Reduce とか上位層で駆逐されてしまうねん
cudaやってるんですけど・・・
カーネル関数起動させるところでエラーが出てしまいます
サンプルコードでアウトなんです
考えられる可能性を挙げていただきたいです
エロい人助けてください
ちなみに、
win7professional32
グラボ１：８６００gs（出力用）
グラボ２：４６０gtx(→cuda)
開発環境：visual studio 2008
質問あればできるものはすべて答えますんでよろしくお願いします
エラーメッセージぐらいのせろやカスが
ttp://developer.amd.com/gpu/ATIStreamSDK/assets/AMD_Evergreen-Family_Instruction_Set_Architecture.pdf
メモリアクセス周りの挙動について言及されているのがなかなか面白い。
コアレス化が余り重要じゃないという話がどういう意味か分かる。
要は、アーキテクチャ的に1スレッドが複数のメモリアクセス命令を同時発行可能で
1wavefront単位で発行された複数のメモリアクセス命令の間だけ
キャッシュ無しアクセスでもキャッシュが有効になっているから
複数のメモリアクセス命令間でコアレス化と同様の効果が得られるらしい。
スイマセン
CUDA error: Kernel execution failed
コンパイルはできていて、
他のマシンでは同じソースコードのプログラムは動かせます
原因は何なんでしょうか
カーネルの起動に失敗したというのだから、起動要件を満たしていないのだろう。
まさかとは思うが、CUDA用のドライバをインストールしていないというオチではあるまいな。
ドライバは入っています
それと、今日起動することに成功しました。
main()の変数宣言のすぐ後に、
cudaSetDevice(1);
を記述したら、それで通りました。
なぜ起動できなかったかは分かりません。
今可能性を探っているのですが、
タイムアウトが起こったのかもしれないと考えています。
CUDAの命令はよく分からんけど、名前から察するに単純に処理するGPU指定してなかっただけじゃ。
8400GSだってCUDA対応なんだし。
デバイスを指定しなければデフォルトのデバイス0、8600GSで実行されるはずだけど。
Capability 2.0以降限定の関数を呼び出しているサンプルコードだったとか？
まあCUDAスレあるし、そっちでやるべきかな。
http://galaxy.u-aizu.ac.jp/trac/note/wiki/Fast_GEMM_Implementation_On_Cypress
４VLIWプロセッサーになるとのことですが、現在のCALでサポートされているsin/cos等の
超越関数命令は、自分で多項式近似計算をしろと言うことなのでしょうか？
今までtレーンが担当してきた命令は
xyzwの複数レーンで1命令を実行するように変更されている。
で、超越関数は3SPを使った1命令で実行される。
53 です。 CORDICやチェビシェフ多項式をWeb上で漁っていたのですが
安心しました。
ソフトウェア無線専用ハード（微妙に矛盾？）でCORDICを使ってるとか話には聞いたことがありますな。
あと自分で実装してみて気づいたんだけど、分岐と言ってもある数を足すか引くかなので、
分岐しないようにしてビット操作に落とせるんですよね。
GPUは同一ワープ内で命令ストリームを共有してますから同じ方向にしか進めない。
Cでいうif-elseは一見分岐だけど、GPUではプレディケート情報によって実行・不実行（あるいは結果に反映させない）を
選択する単一の流れに展開されています。
プレディケート自体はそんなに重たくないです。
むしろ分岐先が増えると増えた分だけ処理時間が増えるだけで。
もう手に入らんの？(´・ω・｀)今日HD5670GETしたのに・・・
http://bluesky23.yu-nagi.com/
うぉ！マジでありがとう
ググっても見つけられなかったんだ(T-T)ｳﾚｼｲ
GPGPU向けに機能追加されたSandy Bridge GPUコアの登場を待ちましょう。
だんごに好かれてしまったからGPGPUも端ッパの技術バリエーションの一つに転落決定だな
ただしアセンブリ言語で書くのは嫌。
これでCUDAできますよね？
俺の夢かなえられますよね？
ひゃっはあああああああああああああああああああ
グレートプログラマー初春
CPU+GPU叩いて作ったほうがはるかにましだがなあ。生産効率が桁違いだわ。
今までDSPやFPGAでやっていた事をCPU+GPUでできる用途ってどんなものがあるの？
考えたけど一つも思い浮かばなかった。
みたいな処理では、GPUでの代替はかなり強烈だよ。
FPGAだとたかだか数百タップの複素FIRフィルタを40～50MHzの動作速度でさばくのにも
現状だと5～15万ぐらいのデバイスがいるし。
大量生産するものなら処理をチップ化して安くあげちゃうんだろうけど、
俺のとこみたいな数のでない無線通信製品だとGPGPUはかなり魅力的。
たぶん、画像検査装置みたいな分野でもGPGPUは強力だと思う。
Windows版は付属していますか？
それどころかCALは(IL含め)2.5で死滅。
かわりにLLVM IR使え。
そんな感じです。
ありがとうございます。
そうですか。。
せっかくIL習得したところなんですが、困りましたね。
ここにはCALのsampleはcalといディレクトリにあるとかいてますが、
旧バージョンから修正されてないだけでしょうか？
無難という意味では、OpenCLですか？
ソフトをやるのかハードをやるのか？
ソフトの場合は上位をやるのか下位をやるのか？
（１）CPU Intel Q9450 (4 posix threads) GCC 4.4.3 （最適化無し、１２５万個／スレッド）
（２）ATI HD4870 ＋ OpenCL (AMD APP SDK 2.4) （最適化無し）
で計算させてみた（行列は正定値実対称の素直な行列）。
ハードはCPU、GPUともに定格で使用。OSはUbuntu x64 10.4 LTS , AMDドライバはCatalyst 11.5
GCC4.4.3とOpenCLで使用したソースコードは略同じものを使用（相違点は OpenCL側コードに__global 指定が付いた程度）
時間測定はC言語側の計算ルーチン呼び出し元でgettimeofday()を使用してマイクロ秒単位で測定。
------------------------------------------------------------------
Posix thread creation : 0.000124(sec)
Thread #1 Exection time: 5.992(sec)
Thread #2 Exection time: 6.08448(sec)
Thread #3 Exection time: 5.9132(sec)
Thread #4 Exection time: 5.91843(sec)
Total Exection time : 6.08452(sec) <ーースレッド中の最大値 + α
------------------------------------------------------------------
GPU Kernel Compile : 1.6e-05(sec)
GPU Kernel Build : 5.02891(sec)
GPU Kernel Creation : 6e-06(sec)
GPU Kernel Set Args. : 2e-06(sec)
GPU Kernel Execution : 4e-05(sec) --- clEnqueueNDRangeKernel ()を挟むgettimeofday()の時間差
Memory mapping(READ MODE) : 5.38053(sec)
＜この間でデータ読み出し＞
Memory UnMapping(from READ MODE) : 0.020504(sec)
OpenCLソースのビルド＆結果データの読み出しまで含めるとGPUが1.7倍遅いが計算実行時間の単純比較だと
6.08452 / 4.0E-5 = 1.5E5 = 15万倍速い！ と言う結果になりました。 いくら何でも１５万倍は速すぎのような気が・・・・(^_^;;)
以上。
ごめんなさい。 ミススペルしてました。
(1) Q9450 4スレッド
------------------------------------------------------------------
Posix thread creation : 0.000124(sec)
Thread #1 Execution time: 5.992(sec)
Thread #2 Execution time: 6.08448(sec)
Thread #3 Execution time: 5.9132(sec)
Thread #4 Execution time: 5.91843(sec)
Total Execution time : 6.08452(sec) <ーースレッド中の最大値 + α
(sizeof float)*3*3*5000000≒180[MB]
これがシステムメモリとVRAM間で往復するから360[MB]
所要時間が2[ms]だから、1[s]に180[GB]も動いてることになる
何か変だ
使った行列は
2.000000E+00, 1.000000E+00, -1.000000E+00
1.000000E+00, 3.000000E+00, 2.000000E+00
-1.000000E+00, 2.000000E+00, 4.000000E+00
ただしデータは対称性の為、(2.000000E+00, 3.000000E+00, 4.000000E+00, 1.000000E+00, 2.000000E+00, -1.000000E+00)
の６成分のみで、 システムメモリ～VRAM間の転送量は
３×３行列 sizeof(cl_float)*6*5000000 = 114MB
固有値 sizeof(cl_float)*3*5000000 = 57MB
固有ベクトル sizeof(cl_float)*9*5000000 = 171MB
反復解法の収束回数 sizeof(cl_int )*5000000 = 19MB
CPUでの解は
Eigen value e1 = 2.855207E-01, Eigen vector1 = ( 6.345872E-01, -5.961551E-01, 4.918314E-01)
Eigen value e2 = 3.143277E+00, Eigen vector2 = ( 7.708385E-01, 5.341269E-01, -3.471549E-01)
Eigen value e3 = 5.571202E+00, Eigen vector3 = ( 5.574225E-02, -5.994226E-01, -7.984895E-01)
<e1,e2> = 5.960464E-08 ← 固有ベクトル間の内積での直交性チェック
<e2,e3> = 0.000000E+00
<e3,e1> = 0.000000E+00
GPUでの解は
Eigen value e1 = 2.855215E-01, Eigen vector1 = ( 6.345873E-01, -5.961550E-01, 4.918314E-01)
Eigen value e2 = 3.143277E+00, Eigen vector2 = ( 7.708384E-01, 5.341271E-01, -3.471551E-01)
Eigen value e3 = 5.571201E+00, Eigen vector3 = ( 5.574221E-02, -5.994227E-01, -7.984894E-01)
<e1,e2> = -4.470348E-08
<e2,e3> = -5.960464E-08
<e3,e1> = 0.000000E+00
で略一致してます。
89 です。
同じ問題を maxima の eigens_by_jacobi で解くと
(%i1) A:matrix([2,1,-1],[1,3,2],[-1,2,4]);
(%i2) eigens_by_jacobi(A);
(%o2) [[0.28552125561224, 3.143277321839643, 5.571201422548121],
[ 0.63458730239817 0.77083835004074 - 0.055742207899264 ]
[ ]
[ - 0.59615502320039 0.53412697029887 0.59942269552653 ]]
[ ]
[ 0.49183141821965 - 0.347155034107 0.79848934767235 ]
（こちらは、固有ベクトルの成分が縦方向に並んでいます）
89です。
rtn = clBuildProgram ( pgm, the_number_of_devices, devices, "-cl-opt-disable", NULL, NULL );
でBUILDしています。
KHRONOSのPDFマニュアル p115に
-cl-opt-disable
This option disables all optimizations. The default is optimizations are enabled.
と記述があります。
またkernel 実行は
rtn = clEnqueueNDRangeKernel ( CommandQueue,
kernel,
1,
NULL,
&pe_size, // 5000000
&group_size, // 64
0,
NULL,
NULL // No triger event will be used.
);
今気がついたのですが、ｐ１３２に
clEnqueueNDRangeKernel returns CL_SUCCESS if the kernel execution was successfully ～～ queued ～～.
Otherwise, it returns one of the following errors:
とありました。 "Execution Time" と上で書いた時間は実行キューへの登録時間でした。
お騒がせしてすみませんでした。
cl_event event;
rtn = clEnqueueNDRangeKernel ( CommandQueue, kernel, 1, NULL, &pe_size /* 5000000 */, &group_size /* 64 */,
0, NULL,
&event <- イベント追加
);
if( event ){
(void)clWaitForEvents( 1, &event );
(void)clGetEventProfilingInfo ( event, CL_PROFILING_COMMAND_START, sizeof(cl_ulong), &device_time_counter[0], NULL );
(void)clGetEventProfilingInfo ( event, CL_PROFILING_COMMAND_END, sizeof(cl_ulong), &device_time_counter[1], NULL );
(void)clReleaseEvent( event );
}
実行時間 device_time_counter[1] - device_time_counter[0] (nsec);
GPU Kernel Compile : 1.5e-05(sec)
GPU Kernel Build : 5.02459(sec)
GPU Kernel Creation : 6e-06(sec)
GPU Kernel Set Args. : 1e-06(sec)
*GPU Kernel Execution : 0.114895(sec)
*C[114MB] memory mapping(READ MODE): 0.0358828(SEC) 3177.01(MB/sec)
*E[ 57MB] memory mapping(READ MODE): 0.0179288(SEC) 3179.24(MB/sec)
*V[171MB] memory mapping(READ MODE): 0.0537894(SEC) 3179.07(MB/sec)
*iter[19MB] memory mapping(READ MODE): 0.00600078(SEC) 3166.26(MB/sec)
*はOpenCLのプロファイリング機能で測定した時間。 それ以外はgettimeofday()を使用して呼び出し元から測定した時間。
結局 6.08452 / 0.114895 = 52.96倍 次期 HD7000 が楽しみになってきました (^_^)。
TOP カテ一覧 スレ一覧 2ch元 削除依頼 ▲
「Cでプログラミングするには人生は短すぎる」か？ (314)
エスパーが質問に答えるスレ (309)
【bzr】Bazaarでバージョン管理 Rev 3 (958)
【GUI】wxWidgets(旧wxWindows) その5【サイザー】 (496)
画像処理 その13 (876)
文字コードの種類は何故複数あるのでしょうか？ (342)
--log9.info------------------
m.c.AT富樫明生を語りたい奴等の広場Part5 (262)
中目黒 ◯○solfa○◯ OPEN (222)
●大阪▲あんだーぐらうんど■ (640)
eggman 【渋谷】 (618)
渋谷道玄坂GASPANIC (811)
◆◆BMS界のTRANCE総合スレ◆◆ (458)
新宿によいクラブありますか？ (201)
■HERBERT■ハーバート総合スレv1.9■HERBERT■ (815)
【イベント】石川のクラブスレ【いきまっし 】 (356)
【詐欺師】YAHOOオークション pt4【撲滅】 (524)
福岡のハウス№1DJって誰よ？VOL3 (918)
【大阪】お も ろ い イベントない？【暇】 (320)
自分が格闘家だったら入場曲は？（クラブ編） (598)
＝＝ASTRIX＝＝ (300)
〓元気ロケッツ〓♂→⇒ (351)
クスリやってる・やってそうなDJ (307)
--log55.com------------------
DDRプレイヤースレ33
beatmaniaIIDX 十段スレ☆232
beatmaniaIIDX 七段スレ ☆47
beatmaniaIIDX 中伝スレ☆41
GITADORA NEXTAGE総合スレ3 【ギタドラ】
jubeat総合スレ 409th Tune
SOUND VOLTEX VIVID WAVE TRACK365
ACポップンミュージック☆2115 ラクリマLV.49