【GPGPU】くだすれCUDAスレ part5【NVIDIA】

おまけ２

韓国は現在でも使っているから異常だ！とか言いそうだが、韓国人に嫁いだ日本人女性の
ブログは噂話程度の内容だし、その他は30年以上前の話ばかり。
じゃあ、日本ではどうなの？と言うと、似たような話があったりする。
「落し紙以前」（斉藤たま、論創社、2005）には、東北の一地方では便器の内側についた
カスを虫歯の穴につめる民間治療があったり、病気の子供を治すために牛糞でくるむ話が
あったり、「トンスル」同様の方法でとった「糞清」？を病気の治療に使う話などが
紹介されている。
まあ、そもそも日本でも韓国でも、排泄物を「汚い」とみなすようになったのは歴史的には
割と最近だったりするわけで・・・。*12

mujinさん流に言うなら、お前の尻も赤いぞ、ってこと。*13




*1：日本と韓国の間の旅行客数の推移　http://www2.ttcn.ne.jp/honkawa/7100.html

*2：http://blog.goo.ne.jp/mpac/e/4aba204bd01800448deb3f509a492d75

*3：人糞説に拘るバカは後を絶たないが・・・。最後には人糞じゃないということを証明しろとか言い出す始末・・・。
*4：2004年7月　http://www.ntv.co.jp/wrs/renewal/ranking/20040704/01.html
*5：ひょっとしたら、逆にテレビの内容を改ざんしてネットに広めた
*6：http://ja.wikipedia.org/wiki/%E3%83%9B%E3%83%B3%E3%82%AA%E3%83%95%E3%82%A7
*7：http://hanja.naver.com/search?query=%EB%98%A5
*8：http://www.gnnews.co.kr/index.html?section=KNJK&flag=detail&code=173716&cate1=KNJ&cate2=KNJK
*9：http://mayanagi.hum.ibaraki.ac.jp/LecRep/04/TextInterp2.html
*10：2巻
*11：http://www.naturalharmony.co.jp/trust/koramu/3sarada-kinenbi.html
*12：人糞を肥料として使っていた日本では、人糞の取り合いで大喧嘩になったりしてる。
*13：http://d.hatena.ne.jp/mujin/20090515/p1

　慶尚北道金泉で生まれ育った小説家、金衍洙（キム・ヨンス）氏は、25歳のときガンギエイを初めて見た。
ソウル・仁寺洞の居酒屋で先輩の詩人に勧められるままに、ピンク色の魚を一切れ、何気なく食べてみた。
その瞬間、彼は「全校生が使用する便所を丸ごと飲み込んだような感じ」に圧倒された。
アンモニア臭が口から3日間消えなかった。
金氏は「人間はなぜ、このような食べ物を食べなければならないのか」と考えた。
人生の中でそれは「死ぬと分かっていながら死に向かって疾走する欲望」だった。
そのため、金氏はガンギエイを「大人の食べ物」と呼んだ。

　ガンギエイを発酵させる過程を見ると、金衍洙氏が便所を連想したのも無理はない。
つぼに石を入れて、その上にワラを敷き、ガンギエイを置いて、繰り返して重ねていく。
そしてつぼの口をしっかり閉めて、暗い倉庫に置いておく。実際には腐らせるのと変わらない。
5、6日後につぼを開けると、息が詰まるほどのアンモニア臭があふれ出てくる。
昔、寒くてガンギエイがあまり発酵しない冬には、堆肥を腐らせる堆肥置き場にガンギエイを置いておくこともあったという。

　それでも、ガンギエイが腐らずに発酵するのは、体内に多くの尿素があるからだ。
ガンギエイを発酵させると尿素が分解されてアンモニアが発生する。
タンパク質が腐敗してアンモニアを発生させるのとは異なる。アンモニアは臭いがきつい毒性物質だが、
ガンギエイから発生するアンモニアは体に害を及ぼすレベルではない。
むしろ、ガンギエイを強アルカリ性にし、腐敗細菌や食中毒細菌の繁殖を抑える。
キムチやチーズのような発酵食品と同様に、ガンギエイの味を覚えると、病みつきになる。
口の中や舌の皮がむけるほどピリッとする味に魅せられてしまう。

（つづく）
http://www.chosunonline.com/site/data/html_dir/2012/08/18/2012081800546.html
http://www.chosunonline.com/site/data/img_dir/2012/08/18/2012081800530_0.jpg

韓国の新聞社がちゃんと証言してるんですけど？


#昔、寒くてガンギエイがあまり発酵しない冬には、堆肥を腐らせる堆肥置き場にガンギエイを置いておくこともあったという。

>>904
君は日本語がちゃんと読めるの？
>>896が何を書いてあるか全く読めないの？
あなた本当に日本人？
エイはサメと同じようにアンモニアを多く含んでいて
腐りにくいために日本の海から離れている山間部で
サメとかがたべられていたのもしらないの？
http://ja.wikipedia.org/wiki/%E3%82%B5%E3%83%A1#.E9.A3.9F.E6.9D.90.E3.81.A8.E3.81.97.E3.81.A6.E3.81.AE.E3.82.B5.E3.83.A1
http://ja.wikipedia.org/wiki/%E3%82%A8%E3%82%A4#.E5.88.A9.E5.AE.B3

韓流依存の低いテレ朝でも韓国ドラマなくなったら
倒産するらしい
他も推して知るべしといったところだな
テレビ局総倒れということだな


422+1 ：名無しさん＠お腹いっぱい。 [sage] ：2012/08/18(土) 19:48:49.96
>>421
そんな事をやったら体力の無いテレ朝は潰れますからｗ
テレ朝の親玉・朝日新聞だってTBSと組んで韓流やってるのにｗ
新聞業界も将来どうなるか分からないし朝日新聞も潰れない保証は無い



もう、命綱といわれる韓国ドラマを一切やめて
一度全局倒産したうえですべてリセットした方がいいわ
批判しながら韓国ドラマはネット配信に取られたくないとかいう
ダブルスタンダードなんか甘いの一言なんだよ
もう、何もかもうんざりだ

テレビ局の給料なんて
持ち株会社化する前のフジテレビの
全社員の給料分でさえ
一日２０時間平均で２００万円台程度なんだから
こんなはした金じゃ給料カットしても
在日だらけの日本のテレビドラマの
放映権ぶんの足しにもならない

だいたい芸能人のギャラが高すぎるんだよ
テレビ局員の給料も芸能人のギャラもすべてリセットだ
　　　　　　,,,
(　ﾟдﾟ)つ┃

君たちはこのスレが何のスレかも分からないの？

なんのスレか以前に
日本最大の右翼掲示板ですから

ネトウヨの陰謀論掲示板

ひとつのスレッドがメモリアクセスする命令を複数回連続して発行したとします。
（それぞれのメモリアドレスは別々です）

この場合、一つ目のメモリアクセス命令が完了するを待つ間に、
同じスレッドの次のメモリアクセス命令が実行を開始するなんてことはあるのでしょうか。

それとも、先のメモリアクセス命令が完了するまで、
次のメモリアクセス命令は実行されずにスレッド自体が止まるのでしょうか。

この辺りの挙動はグローバルメモリへのアクセスでも、
シェアードメモリへのアクセスでも同じでしょうか。

>>910
シェアードだろうがグローバルだろうが、両方ともあり得る

>>911
ということは、シェアードメモリにアクセスする命令を複数回連続して発行した時も、
バンク・コンフリクトは起こり得るのですか。

つまり最適化のためには、並列に実行される複数のスレッドからのアクセスだけでなく、
ひとつのスレッドからアクセスするシェアードメモリについても、
バンク・コンフリクトを意識しないといけないのですね。

まぁ要するに韓国が絡むとどこでも気持ち悪くなるという話だな。

ｽﾚ違いの話をいつまでも続けるんじゃねぇよカスが。

ネトウヨの妄想陰謀論は他でやれ

このスレを嫌韓レスで荒らしまわる
「情報ライブ　ミヤネ屋」に出ている
川田裕美とかいう読売テレビの
馬鹿女狐アナウンサーに抗議だ



川田の女狐の上司のtwitter
http://twitter.com/TakeshiMori

川田の女狐のtwitter
http://twitter.com/hiromikawata

読売テレビに直接抗議する場合
http://www.ytv.co.jp/bangumishinsa/opinion.php

>>912
読みまちがえた、1つのスレッドが連続してメモリを読み書きする場合か。
それならシリアルに実行される

>>916
ありがとうございます、了解しました。

北森瓦版 - “Hyper-Q”によりHPC性能を高めるNVIDIAのGK110
http://northwood.blog60.fc2.com/blog-entry-6178.html

>>918
う〜ん、よー分からんｗ

>>918
うむ、よくわからんｗ

CPU→GPUのデータロードのボトルネック解消(?)のように思えるけど、
ボトルネックになってるのは大抵GPU内の計算かVRAM←→GPUのアクセスじゃなかったっけ？

つまりOut-Of-Orderのメモリ転送ってことなのかな？
QってQueueを意味してたり。

1本当たりの帯域はへるけど最大３２本までメモリ転送が平行でできるってこと
のように見える。データ待ちを減らしてMPIのような全体協調アプリの性能を上げる
というようなものかな

ひょっとして律速がGPU内の計算かVRAMアクセスかとかいうレベルじゃなくて
ばんばん主記憶からデータを持ってこないといけないようなGPGPUプログラムに対する
向上策って感じなのかな？

例えばDSPの代用みたいにリアルタイム処理で使うなら頻繁にCPU←→GPUがあるだろうから効果あるのかも。

いままで、非同期処理にストリームを使っていたところ、
結局ジョブのキューが1個しかないからいろいろ制限がかかってた。
それを、キューをコア数分だけ用意することによって･･･
という話かと考えた

Geforce580からGeforce690に変更し，
SDKのサンプルを動かしたのですが，動作が遅いです.
そもそもCPUの動作も以前より格段に悪いです.
なぜかわかる方いたら教えていただけないでしょうか？

ちなみに
ドライバーやその他の設定は580のときのままです.
また，cuda4.0を使用しOSはwin7 64bitです.

>>927
CUDAの性能に限って言えば580が一番高いんじゃないの?
6xxシリーズは単精度の浮動小数点演算以外の性能は削りまくってる
みたいだし。命令ごとのクロック数の表をみたら一目瞭然でしょ。
自分は580 SLIだけど、Kepler 2が一般消費者向けに販売されるまで様子見してる。

これ、参考までに貼っておくね。
「命令別スループット」の表を見てちょうだい。

GTX680のグラフィック・GPGPU性能を調べる
http://dokumaru.wordpress.com/2012/03/27/gtx680-spec/

>>929
マジかよ… なんじゃこりゃ…


I32加算：8/7、1、1
論理演算：24/17、1、1

>>930
この話、今年の3月頃に2chで出てたよ。
http://anago.2ch.net/test/read.cgi/software/1311428038/633-643n

>>929
うわぁ･･･。
ここまでヒドいとは･･･。

よく見たら680/690のGPGPU性能が芳しくないという話はこのスレでも出ているな。
>>635には同じリンクが紹介されてる。

これ大事だから次スレのテンプレに入れておいたほうがいいな。

Kepler2だとこれ改善されるの？
↓
I32加算：8/7、1、1
論理演算：24/17、1、1

>>931
っつか、団子さんがCUDA関連のスレにいることに驚いたｗ

>>935
グラフィックに不要な部分をばっさり切ったのがkepler1なので当然改善するでしょう。
問題はマス向けのお値段でkepler2が出るのかどうか心配な点。。

チョンと在日チョンは直ぐにファビョります(火病)。
これは劣等遺伝子を持つチョン特有のもので、最近もジャップスレ立てていた犬猿単細胞がファビョったので
容易くレイプ劣等チョン遺伝子を持っている事が分かります。分かりやすいですね？

チョン語と日本語はそっくりな発音が多いとされる。これは、チョンが朝鮮併合後に日本語から
取り入れたものである。つまり、当時のチョン語では単語が足りず、文法も原始的だったのだ。
併合後、日本から取り入れたために発音その他が同じものが多く残っているのである。
当時のチョンは、当時の日本人がびっくりするくらいの土人乞食そのものであったのだ。

日本の朝鮮併合前の写真である(当時のチョン半島の写真)。
当時のチョン半島が如何に土人乞食の文化がまるでないカス民族であったかが分かるだろう。
これがチョン歴史の真実である。乳丸出しの民族衣装がチョンのリアルの民族衣装である。

ttp://www13.plala.or.jp/keibagogo/img/chosen2.jpg
ttp://nandakorea.sakura.ne.jp/img/hitob.jpg
ttp://www13.plala.or.jp/keibagogo/img/chosen3.jpg
ttp://blog-imgs-32.fc2.com/x/i/a/xianxian8181/CHONONNA.jpg

なのでクソを喰う風習も真実ならば、嫁入り前の娘を父親がパンパンと犯すレイプ文化も
リアルにチョンの実態そのものである。チョンが歴史を捏造するのは、こういう恥辱に
まみれた人以下の文化しかない真実を隠そうとしているだけなのだが、現在もチョンは
犯罪とレイプしかしない屑以下の存在である。何でもかんでも捏造し、まさに生きている
価値すらない。チョンと在日チョンは一匹残らず皆殺しにしなければならない。

数学五輪に各国が力を入れていないだけの話ですね。その間隙をぬって一位になったチョンが
工業力で優秀ですか？ノーベル賞を一人でも輩出しましたか？
製品をコピーして低価格で販売するだけしか能がチョンが何を言っているのでしょうか。
自己満足は楽しいですか？

昨日の続きです.
GTX690に変更したところ，580で動作していたプログラムが
動作しない(途中で落ちるor動作が鈍い)くなったんですが
なぜだかわかりますか??

CUDA4.0とGTX690の相性が悪いのでしょうか？？

CUDA 4.0で普通に動いていたプログラムが
CUDA 4.2にしただけで速度が1/5ぐらいになったぞ。
ひょっとしてCUDA 4.2って地雷?

俺も4.0から4.2にした時に、
コンパイル時のメッセージで
「とりあえず浮動小数点は全部倍精度にしとくね！」
みたいなメッセージが出てた気がするけど
単精度メインの人だと処理速度に影響出たりするのかな？
あんまり検証してなくてスマン

ほとんど整数演算だから一般的な使い方とは違うけど、それにしても遅すぎ!
4.0に戻して検証してみるよ。

遅くなったと思ったのは勘違いだった。スマソ。
でも4.0で何ら問題なくコンパイルできたプログラムが
4.2だとReleaseにするとコンパイラが落ちてビルドできなかったりと
挙動が不審すぎる。4.1でも4.2と変わらない。

…よくみたらマクロ使いまくった変態的プログラミングをしていた箇所で
ciccが落ちてた。この箇所書きなおさないといけないなorz

939の問題なのですが解決しました。
ただの接触不良みたいでした.

もうひとつ質問なのですが，
GPUに重い処理をさせたらOSごと落ちることありますよね？？
あれってなぜなるのですが，また，GPUの性能とどのようにかかわりがあるのですか？

すみません。よろしくお願いします。

熱排気がダメで高温落ちじゃないの

電力オーバーとか

なるほど...
分かりました！！　調べてみます。
ありがとうございました。

やっぱりまだまだGPUプログラミングは厄介だね･･･

>>944
windows7だとGPU処理が一定時間返ってこないとディスプレイ落ちるが..
それとは違うんだな？

>>947
Windowsはドライバの応答がなくなると固まる
10秒以上かかりそうな処理は適当な間隔で戻してループしないとダメなようだ

その制限、解除出来るけどね

>>951
土下座しますので教えてください｡

>>951
できねーよ
できるならここに詳細を書いてみろ(プゲラ

>>947
>>950-953
制限の解除ってこれのことだろ。

CUDAカーネル実行のタイムアウト - PukiWiki Plus!
http://imd.naist.jp/~fujis/cgi-bin/wiki/index.php?CUDA%A5%AB%A1%BC%A5%CD%A5%EB%BC%C2%B9%D4%A4%CE%A5%BF%A5%A4%A5%E0%A5%A2%A5%A6%A5%C8

これもだいぶ前に2chで出てた話。「プゲラ」じゃねーよｗ

nVIDIAが出してるツール使えばレジストリエディタさえ使わずに制限解除できる。

　　　　　　ｸ　　　 ｸ　｜|　プ 　//
　　　　　　ｽ　 ｸ　ｽ　　| | │　//
　　　　　　 ／ ｽ　　　　| | ッ //　 　ｸ　　　ｸ　 ｜|. プ 　//
　　　　　　 ／　　　　　　 　 //　　　ｽ　ｸ　ｽ ＿ | | │　//
　　　　　　　　 ／￣￣＼　　　　　／　 ｽ　　 ─ | | ッ //
　　　　　　　／　　_ノ　　.＼　　 　 ／　　　　　　　　　//
　　　　　　 ｜　　（ ＞）（＜）　 　 　　　＿＿＿_
.　　　 　 　 ｜　　⌒（__人__） 　 　　.／　⌒　 ⌒＼
　　　　　　　 |　　　　｀ Y⌒l　　　 ／　 （＞）　（＜）＼
.　　　　 　 　 |　　　　. 人__ ヽ　／ 　::::::⌒（__人__）⌒ ＼
　　　　　　 　ヽ 　 　 　 　 }|　| | 　 　 　　 ｀ Y⌒ l__　　　|
　　　　　　　　 ヽ　　　　ノ､|　|　＼ 　 　　　 人＿　ヽ ／
. 　 　 　 　 ／^l 　　　　　 / /　 　,─l　　　　　　 ヽ　＼

竹島問題で　韓ドラとＫ−ＰＯＰが消える！？

　竹島問題を巡る日本と韓国の対立はテレビ界でも大きな影響が出始めている。韓流ドラマとＫ―ＰＯＰが近いうちに
消滅するのではと囁かれているのだ。

　この３年でテレビで放送された韓流ドラマは約５００番組、日本でデビューしたＫ―ＰＯＰのグループは３０組以上
にも上る。テレビ局にとってはドル箱といっていいコンテンツだった。

「例えば、韓ドラはテレビ局には安く、そこそこ視聴率が稼げる番組として重宝がられました。買い付け価格は一時上がり
ましたが、最近はダウンし、初回購入の一番組の単価が５００万円程度というものもあった。しかも、３回まで再放送が
認められるのでＢＳやＣＳが飛びついていた。一方、Ｋ―ＰＯＰは音楽番組に出るのはプロモーションの一環で出演料が
かからなかったり、レコード会社が立て替えて支払ってくれるケースも多かったのでメリットが大きかった」（事情通）

　だが、例の問題で状況が一変した。すでに日本テレビやテレビ朝日、ＴＢＳなどは年末年始の番組編成から韓ドラ枠を
大削減する方針を固めたという。音楽番組もＫ―ＰＯＰのゲスト出演を極力減らしていく方針だという。また、
韓国びいきと批判されているフジテレビでさえも韓ドラからの撤退、Ｋ―ＰＯＰの露出を減らしているという。

「李明博大統領が竹島上陸、ロンドン五輪でサッカー韓国代表メンバーが『独島は我が領土』というメッセージを掲げ、
韓流スターらが竹島に泳いで渡るという反日的ニュースが流れてから、ＢＳとＣＳで視聴者の接触率が激減しました。
竹島問題で韓国の反日的なパフォーマンスが続く限り、来年４月の番組編成で韓ドラとＫ―ＰＯＰは地上波から消滅する
可能性が高い。ＮＨＫも韓ドラをかなり流していますが、右に倣えだと思います」（民放編成マン）

　テレビ局にとっては韓ドラに代わるコンテンツが見つからないのは頭が痛いが、Ｋ―ＰＯＰに代わるダンスミュージック
を発掘中とか。ＫＡＲＡもチャン・グンソクも来春以降は見ることができなくなる？

http://gendai.net/articles/view/geino/138538

>>959
北チョン芸能ヤクザとべったりともっぱらの
日刊ゲンダイの記事かよ
北チョン芸能ヤクザ自身が書いているともいわれているな

>>957
北チョン芸能ヤクザとべったりともっぱらの
日刊ゲンダイの記事かよ
北チョン芸能ヤクザ自身が書いているともいわれているな

嫌いな国ランキング 2012 投票受付中！
http://japaneserank.enq1.shinobi.jp/vote/101289/

このスレは何故かネトウヨの陰謀論がよく沸くね

num = 0;
for(i = 0;i < 100;i++)
{
if(条件)
{
list[num++] = i;
}
}
のような感じである条件をみたすものだけを
リストに追加、その個数をカウントしたいのですが
CUDAで実装可能でしょうか。

可能は可能。極論言えばそのループを1スレッドで回せばいい

効率よくできるかどうかは詳細部分に依存する。

苦手な部類だな。
CPUの方がよっぽど手っ取り早くやってのけられる。
仮にiの最大値が100なんてしょぼい数じゃなく巨大なら、
スレッド数で分割してスレッドごとにlistを作って後で繋ぐこともできるけど。

リストに順に詰めていくことに意味があるのなら、それはつまり
i==nのときの処理の後でないとi==n+1のときの処理を行えないことを意味するので、
並列化に向いていない。
100じゃなくてもっと大きい数でなおかつ格納順を気にしなくていいなら>>964の言ってるとおり。

あとカウントはAtomicAddとかを使うよりトーナメント方式(？あの隣同士をlogN段階に分けて
足すやつ)のほうが早かった気がする。

ありがとうございます。順序は関係ないので
>>964さんの方法で分割しようと思います。
ただ、listを繋ぐために要素をコピーしていくのも
時間がかかりそうです（自分の理解が浅いかもしれませんが）
CPUが手っ取り早いかもしれません。

条件の真がint の1になるなら分岐せずに
= (条件) * i
にするとか

順序が関係ないなら、いろいろやりようがあるな。
メモリに余裕があれば、条件が真か偽かで1/0を入れるところまで
馬鹿並列で計算してしまって、それからリダクションみたいな操作
（たぶん上で言ってるトーナメント方式ってのと一緒）で詰めるとか

ごめんなさい、リダクションです。
965では正しい呼び方がわからずにイメージでトーナメントって語を使っただけなので
間違って用語を覚えないようトーナメントは忘れてください。


ググってわかったんだけど、リダクションって結合法則さえ満たせば交換法則は
満たしていない演算でも適用可能なんだね。まさに3Dの行列積にぴったりじゃん！
応用する場面は俺には思いつかないけど…

CUDA - Reduction
http://sammaya.garyoutensei.com/gpgpu/reduction.html

>>962
prefix scanを使えばきれいに並列化できる
thrustのcopy_ifを参照

時代はOpenCLじゃないのか?
俺様のみのCUDAって先がないって感じじゃないか

事実上CUDAが標準じゃないの

openclの時代は来るだろうし誰も期待してない訳ではなかろ

ただまだcuda脅かす感じじゃないよな
先はどうなんだろ、まだ見ぬkepler2の動的並列化が標準化される頃にはまた状況変わってるだろうしなあ
xeon phiとか出てくるし

らくちんにプログラムしたけりゃCPUで、性能欲しけりゃCUDAで、って感じでOpenCLはその中間にある。

中間とはいっても、だいぶCUDA側に寄ってるわけで、
CPUプログラミングからわざわざやってくる人は性能を求めてるわけだからOpenCLを
すっ飛ばしてCUDAやるんじゃないかっていう。

しかしコンシューマ用のKeplerを見る限りはなんかnVidiaは安心してサボってるような印象は受ける。
Geforce 7xxシリーズもGPGPUとしては使いにくいものになりそうなので
MaxwellでFermiの上位互換的なものがくることに期待。
でもたぶんこない。コンシューマ向けのGPGPU体験/普及キャンペーンはFermiで終了したっぽい。

>>974
7xxはFermiの「上位互換」だと思ってたけど…
コアはTesla K20と同じGK110でしょ?

>>976
>>975のソースはこれね。

NVIDIA's Monster GPU for Tesla K20, 2013 GeForce and Quadro Cards
http://vr-zone.com/articles/nvidia-s-monster-gpu-for-tesla-k20-2013-geforce-and-quadro-cards/15884.html

性能がほしいからGPU使うのにOpenCL使って性能でないとか馬鹿すぎるので
みんなCUDA使ってるよ

とんがった用途に汎用環境とか、根本的に矛盾してるよね。

Radeon/ATI stream どうなった？

Occupancyを上げるためにPTXでコードを書き直してレジスタ数を削ったり、
48Kしかない共有メモリの使い方をああでもないこうでもないと工夫しているというのに、
NVIDIAのカードでOpenCLを使うなんて考えられない。

CUDAってFORTRANみたいな感じになって、デスクトップのiGPU付きCPUでGPUは
汎用環境用に進化していくんじゃないか。
デスクトップdGPUに重要なのはゲーム、それ以外(GPGPU用)は
あんまり重要ではないから最新のデスクトップ用ケプラーでは
GPGPUを軽視したものにしたんだろ。
今後はとんがった用途にデスクトップ用GPUなんて使うな、
GPGPU用GPUを使えになるんじゃないのか。
デスクトップPC,iGPUの性能向上でゲーム以外の用途では
あんまりVGA要らなくなってきているってのが現状だし

APUはどこへ行くんだろうか

>>975
GK110をそのままスケールダウンしたものがミドルやそれ以下にくればいいけど、
7xxの最上位以外はGK104/106/107のマイナーチェンジなんじゃないかな？
4xx→5xxがそうであったように。
それだとTeslaの名前が変わっただけで俺のような下々の人間にとっちゃ何も変わらんのよ。

寂しいけど>>980の言うようにこうなっちゃいそうなんだよね。
>今後はとんがった用途にデスクトップ用GPUなんて使うな、
>GPGPU用GPUを使えになるんじゃないのか。

>>982
個人的にはGK110が780に来てくれさえすればいいんだけど、
別にローエンドにこないと決まったわけじゃないでしょ。

とんがった用途って具体的に何を使って何をしているんだ?

今発売中のgeforceの中では、どれが一番GPGPUに適してます？
6xxシリーズはダメっぽいって感じでいいのかな？

>>985
580/590の二択しかない。

>>985
でも単精度浮動小数点数だけでいいっていうんなら680/690もあり。
Tesla K10なんてのもあるぐらいだし、需要はあるんだろう。

OpenCLは全てが中途半端
OpenACCで学会の論文とかチラホラ出始めてるけど、当初の予想よりかなり良さそうなので、そのままOpenCLは駆逐される気がする。

>986,987
ふーむ。ありがとう

ここのひとはC++ AMPにも手を出してる？
OpenCLみたいに中途半端なんだろうか。

今はCUDA触れる環境がないから実験で確認できないのだが、
Warpと__syncthreads関数との関係について質問したい。

「はじめてのCUDAプログラミング」という本には、
俺の記憶では次のような意味に取れる事が書かれていたような気がする。

条件分岐にさしかかった場合、
Warp内の「全てのスレッド」がまず一方の分岐を処理し、
本来そちらには行かないはずのスレッドは計算結果を捨てる。
それからもう一方の分岐をまた「全てのスレッド」が通り、
本来そちらには行かないはずのスレッドは計算結果を捨てる。

しかし「CUDA BY EXAMPLE」という本には、これとは違い、
分岐した各スレッドは分岐のもう一方の道は通らないと書かれている。
だから、ifの中に__syncthreadsがあってそこを通らないスレッドがあると、
全てのスレッドが__syncthreadsに到達するという条件が決して満たされず
ハングアップするのだという。

前者の解釈ならifの中に__syncthreadsがあっても問題ない。
最後に計算結果を捨てるが、とりあえず全スレッドがそこを通る。
後者の解釈ならマズイことになる。

本当はどっちの解釈が正解？
それとも実装依存で、環境によってどちらもあり得る？

>>991
if (...) {
....
__syncthreads();
...
} else {
....
__syncthreads();
....
}
は問題ない。昔エミュレータでエラーになることはあったけどGPUでは問題なく動作していた。

>>992
if (...) {
....
__syncthreads();
...
} else {
....
}

でWarp内のスレッドが分かれるの場合の話。

「はじめてのCUDAプログラミング」の解釈だと問題ないと思われるけど、
「CUDA BY EXAMPLE」だとはっきりマズイと書かれている（日本語版73ページ）。

>>991
少なくともWarp内の全てのスレッドの分岐が同方向なら、
そちらには行かないはずでは。
__syncthreadsがWarp単位で数えていたか、
thread単位で数えていたかは分からんが、
少なくとも全Warpが同じ__syncthreadsを
通る必要はあるだろう。

>>992
MPIのbarrierと違って数だけ辻褄合わせても駄目で
ソースの同じ行にある__syncthreadsを通らなきゃいけないとか
制限無かったっけ？

>>991
CUDA4.2のCUDA_C_Programming_Guide.pdfの89ページには
以下の記述がある。
>__syncthreads() is allowed in conditional code but only if the conditional
>evaluates identically across the entire thread block, otherwise the code execution is
>likely to hang or produce unintended side effects.

仕様上は全スレッドで同じ__syncthreadsを実行しろと書いてあるから、
たとえ今のハードやコンパイラでたまたまOKでも
よろしくない書き方だろう。

>>995でFAみたいだね。結局二つの本のどちらも間違ってなかったんじゃないかな？

っていうか>>991をよく見ると「はじめてのCUDAプログラミング」にしたって、
Warp内またはBlock内で分岐が割れるところに_syncthreads()を置いていい、とは書いてないみたいだし。

俺はその本持ってないけど、プログラムの動き方の説明あたりだと予想。
処理時間が片方はT1、片方はT2かかる分岐でなおかつWarp内で割れる場合だと、
処理時間はmax(T1,T2)じゃなくてT1＋T2かかる的な話の。

次スレどうしよう。23:15まで誰も立てないようなら俺が立ててみようかな。

次スレ立てました

【GPGPU】くだすれCUDAスレ part6【NVIDIA】
http://toro.2ch.net/test/read.cgi/tech/1348409867/l50

みんなレスありがと。

>>995
俺も Programming_Guide 読んでみた。
確かに __syncthreads() を通るスレッドと通らないスレッドがあるとマズイと書かれてるね。
仕様なんで、ちゃんと従うようにするよ。

>>996
図書館で「はじめてのCUDAプログラミング」をもう一度確認したら、次のように書かれてた。
ちょっと長いけど引用する（改行は俺が勝手に入れた）。

-----( 52ページから引用 )-------------------------------
if (threadIdx.x & 1)
  a = a + 2;
else
  a = a + 1;
　この例では「スレッド番号」が「奇数」の場合と「偶数」の場合で異なる処理をします。
　どちらの方向に分岐するスレッドも、「ウォープ」内にあります。
このとき、「a=a+2」と「a=a+1」の両方の命令が実行されます。
単純に両方実行すると「a=a+3」と同じことになってしまいますが、
「a=a+2」のほうはスレッド番号が奇数のスレッドのみ有効に、
「a=a+1」のほうはスレッド番号が偶数のスレッドのみ有効にすることで実現されます。
　ここで加算命令を実行した結果は有効なスレッドに対してのみ反映されます。
　有効でないスレッドに関しては、命令を実行していないのと同じ状態になります。
--------------------------------------------------------

この「命令」や「実行」、「反映」、「実行していないのと同じ状態」
という言葉がやや曖昧ではっきりしない印象だけどね。
「a=a+2」の方を __syncthreads() に変えたら偶数番目のスレッドは実行される？
それとも命令を実行していないのと同じ状態になる？
この説明だけではよく分からん。

>>998
「命令を実行していないのと同じ状態」だからつまりそういうことでいいんじゃないかな。
俺が前もって答え（>>995）を知ってるからそう思うのかもしれないけど。

クゥ〜〜･･･　　ッダ！！

このスレッドは１０００を超えました。
もう書けないので、新しいスレッドを立ててくださいです。。。