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

>>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）を知ってるからそう思うのかもしれないけど。

クゥ〜〜･･･　　ッダ！！

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