【GPGPU】くだすれCUDAスレ part7【NVIDIA】 [転載禁止]©2ch.net

**デフォルトの名無しさん** · 2014/11/20(木) 23:14:46.66

このスレッドは、他のスレッドでは書き込めない超低レベル、
もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。
CUDA使いが優しくコメントを返しますが、
お礼はCUDAの布教と初心者の救済をお願いします。

CUDA・HomePage
ttp://developer.nvidia.com/category/zone/cuda-zone

関連スレ
GPGPU#5
ttp://peace.2ch.net/test/read.cgi/tech/1281876470/l50

前スレ
【GPGPU】くだすれCUDAスレ【NVIDIA】
ttp://pc12.2ch.net/test/read.cgi/tech/1206152032/
【GPGPU】くだすれCUDAスレ pert2【NVIDIA】
ttp://pc12.2ch.net/test/read.cgi/tech/1254997777/
【GPGPU】くだすれCUDAスレ pert3【NVIDIA】
ttp://hibari.2ch.net/test/read.cgi/tech/1271587710/
【GPGPU】くだすれCUDAスレ pert4【NVIDIA】
ttp://hibari.2ch.net/test/read.cgi/tech/1291467433/
【GPGPU】くだすれCUDAスレ part5【NVIDIA】
http://toro.2ch.net/test/read.cgi/tech/1314104886/
【GPGPU】くだすれCUDAスレ part6【NVIDIA】
ttp://peace.2ch.net/test/read.cgi/tech/1348409867/

**デフォルトの名無しさん** · 2016/03/04(金) 02:24:38.67

自覚のない荒らし
それが団子の糞風味クオリティ

**デフォルトの名無しさん** · 2016/03/04(金) 07:49:00.48

>>825
何が嘘なのか具体的に示せよ
俺は具体的に示したろ
糞団子は去れ

**デフォルトの名無しさん** · 2016/03/04(金) 07:50:34.02

>>826
＞なぜならキャッシュはDRAMの読み書きを高速化するためにあるから
そんなパソコン教室の初日に教わる様な内容を語るスレじゃねぇよ
糞団子は去れ！

**デフォルトの名無しさん** · 2016/03/04(金) 08:15:25.08

匿名で俺って言い張られてもｗ

**デフォルトの名無しさん** · 2016/03/04(金) 09:40:06.12

>>830
流れが読めないなら黙ってればいいのに

**デフォルトの名無しさん** · 2016/03/04(金) 09:48:28.02

Talk me.

**デフォルトの名無しさん** · 2016/03/04(金) 10:12:42.01

オレオレ、俺だよ俺
俺だよ

かまってよ

ってかｗ

**デフォルトの名無しさん** · 2016/03/04(金) 11:01:16.11

お前つまんねぇよ

**デフォルトの名無しさん** · 2016/03/04(金) 11:46:47.03

団子ちゃん悔しいだろホントはおともだちをつくってワイワイ話したいんだろ？
じゃあ自作板帰ろうな

**デフォルトの名無しさん** · 2016/03/04(金) 12:02:03.20

オレオレ主張したいんならコテハンつけろよ

**デフォルトの名無しさん** · 2016/03/04(金) 15:46:55.78

俺だよ

**デフォルトの名無しさん** · 2016/03/04(金) 19:33:33.59

答えに窮すると何故か匿名になる
安定の自演糞団子

**デフォルトの名無しさん** · 2016/03/04(金) 19:39:28.79

自分がやってるからって
他人もそうだと思うなよ
病気の人

**デフォルトの名無しさん** · 2016/03/04(金) 21:27:27.28

まだやってんのかよっw。
シェアードメモリはローカルメモリだよwww。

**,,・´∀｀・,,）っ-○○○** · 2016/03/04(金) 21:40:33.21

方言のない一般名称でいうと「スクラッチパッドメモリ」か

**デフォルトの名無しさん** · 2016/03/04(金) 21:49:33.06

ぼろ出してもコテでい続けるのは正直苦痛だろ？

な、団子

**デフォルトの名無しさん** · 2016/03/04(金) 22:01:24.79

>>840
>>818

**,,・´∀｀・,,）っ-○○○** · 2016/03/04(金) 23:13:53.73

名無しばっかりだと不安で団子認定に必死な病気の癖に

**デフォルトの名無しさん** · 2016/03/05(土) 04:38:45.77

いい加減通常時に戻ってください
正直どっちもどっちなんだよね
どっちも邪魔

**デフォルトの名無しさん** · 2016/03/05(土) 07:19:08.49

マジで糞団子は去ってくれ
嫌われてるの自覚しろ

**デフォルトの名無しさん** · 2016/03/05(土) 07:44:33.94

コテの人はいい加減ブログでもなんでも始めたらいいよ
荒らし行為だよ

コテ煽る人もやめてよ

**デフォルトの名無しさん** · 2016/03/05(土) 09:30:50.56

団子さん
お願いです。
あなたが来てからこのスレは機能しなくなりました。
ここは数少ないCUDAの情報交換の場です。
もう、ここに来ないで下さい。

**デフォルトの名無しさん** · 2016/03/05(土) 09:48:51.08

実際の荒しは倍精度、精神病の人ですけどね

**デフォルトの名無しさん** · 2016/03/05(土) 10:46:20.52

こいつ一人が団子団子連呼してて
すげー気持ちわりーよ

**デフォルトの名無しさん** · 2016/03/05(土) 10:51:01.60

好きなんだろ
付き合っちゃいなよ

**デフォルトの名無しさん** · 2016/03/05(土) 11:18:09.46

最近は傷心の団子をおもちゃにするうんこごっこが流行ってるのか？

**デフォルトの名無しさん** · 2016/03/05(土) 11:20:03.53

いいえ
精神病の人をどうバカにするかが流行ってます

**デフォルトの名無しさん** · 2016/03/05(土) 14:14:59.66

つーかまともな話しようとすると団子認定されるのは書き込む気を無くさせるためなのだろうか。
技術的な話をすると内容がどうでも団子にされるから、そういう話が理解できないのと、ある意味団子崇拝が酷いのだろうな。

**デフォルトの名無しさん** · 2016/03/05(土) 15:25:30.62

>>849
団子はある程度(正しいかは別にして)技術的話が出来るけど、
その人はAMDは倍精度がすごいよ、GPGPUの勝者はAMDで決まり！と
I　love 団子だけだからな。

**デフォルトの名無しさん** · 2016/03/05(土) 16:26:59.56

あんまり団子をおもちゃにするなって
ウンコの匂いが充満するじゃないかｗ

**デフォルトの名無しさん** · 2016/03/05(土) 16:43:06.41

こんにちは。

CUDA では、__device__ 修飾された関数は、GPU内部から呼び出せますが、
再帰コールが出来ないのはなぜですか？

以下はできますよね？

__device__ int aaa(int x) {
　return x * 2;
}

__device__ int bbb(int y) {
　return aaa(y) + 3;
}

では、以下は駄目ですか？
__device__ int kaijou(int y) {
　if ( y == 1 ) {
　　return y;
　}
　else {
　　return y * kaijou(y-1);
　}
}

**デフォルトの名無しさん** · 2016/03/05(土) 17:05:18.01

団子は答えなくて結構

**デフォルトの名無しさん** · 2016/03/05(土) 17:17:05.46

ハード的に絶対できないと言うことも無いのだろうが、スタックどれだけ使うか計算できないってのが大きいんじゃね。

**デフォルトの名無しさん** · 2016/03/05(土) 19:01:23.69

最適化の邪魔になるから最初から再帰はやめてねーんってことかな。

**デフォルトの名無しさん** · 2016/03/05(土) 20:36:21.36

>>858
おまえが正解答えてやれば
あれずに済むだろ

**,,・´∀｀・,,）っ-○○○** · 2016/03/05(土) 21:10:49.87

>>861
俺ならその無駄レス書いてる暇で書けるぞ>>858は無能だから理解できないだけで

>>859
CUDAの自称コアが要素毎プレディケートつきのSIMDユニットを
マルチコアって言い張ってるだけから本物の独立したコアでないから
で終わりなんだけどね。
命令ポインタは「自称コア」ごとに独立してないんだ

「分岐」も本物の分岐ではなく、if-else節を全部一列のコードに展開してプレディケート
しているだけ。プレディケートビットの全要素が0あるいは1になって初めて関数を
飛び越えることができる。

SIMDの要素（自称）ごとに命令ポインタと関数コールスタックを管理できて別々の
命令ステートを独立して持てるならそれは既にSIMDじゃない

**,,・´∀｀・,,）っ-○○○** · 2016/03/05(土) 21:25:15.51

ベクトルを隠蔽しない、1ワープを1つのスレッドとして処理できるプログラミングモデルが
提供されるなら再帰は理屈上可能
それはCPUのSIMDの在り方と同じってことだけど

**デフォルトの名無しさん** · 2016/03/05(土) 21:53:08.36

独立したコアという前提だから、コア毎のレジスタで表現できる範囲の処理しか出来ないって事だと思ってたけど違うのか。

**デフォルトの名無しさん** · 2016/03/05(土) 22:00:09.02

基本はSIMDプロセッサでの本当の分岐だよ。
簡単な処理ならpredicate実行で済ますけどこれはモダンなプロセッサではSIMT、SIMD、に限らず一般的な手法

状態保存のメモリを最小化するためにdivergent pathを深さ優先でスキャンしていくから、どこかのdivergent pathで無限ループになるとwarp内の別threadには何時までたっても制御が帰ってこないけどね。

**,,・´∀｀・,,）っ-○○○** · 2016/03/05(土) 22:00:10.11

再帰だけならCell SPUすら出来てるからね
命令ポインタの制約のある偽りのコアであること以上の理由はないでしょう

末尾再帰をループに展開するオプションすら用意してないのは処理系の実装の手抜きだろうと思うけど

**デフォルトの名無しさん** · 2016/03/05(土) 22:04:28.99

可能な限り手を抜くのがハードにとってのＧＰＧＰＵという物かも

**,,・´∀｀・,,）っ-○○○** · 2016/03/05(土) 22:12:52.44

実害は具体的にはバックトラッキング型の探索アルゴリズムが使えない
（実装方法に制限がある）とかかな

**デフォルトの名無しさん** · 2016/03/05(土) 22:24:17.32

cudaは性能出すために最適化しだすと、結局SIMD構造を意識して記述することになる。
で結局、SIMTでプログラミングするメリットを感じなくなる。

scatter,gatherも便利なのだけど、それしかないから、性能出すためにアラインメント気にしながら連続アドレスにアクセスするように記述していると、普通のSIMDのようにアドレス１つ計算すればすむのに無駄だなぁと。

で、複数のアドレスを受け取ったハードウェアが同じキャッシュラインに乗っていることを検出してメモリアクセス回数を最小化するとか、バカじゃねと。

まあGCNはそこら辺の無駄に対応していて美しいのだが、いかんせんソフトやサポート体制がクソ過ぎて。

**デフォルトの名無しさん** · 2016/03/05(土) 23:05:45.56

>>865
>状態保存のメモリを最小化するためにdivergent pathを深さ優先でスキャンしていく
横からですが参考になります。
再帰がまずいっていうより結果として条件分岐をたくさん通るのがまずいって感じですかね。

**デフォルトの名無しさん** · 2016/03/06(日) 08:40:29.75

最近は、5,000円くらいの低価格GPUでも、Keplerコアを使っていて、
Fermiコアの上位互換だから、再帰コールも出来るのですか？

でも、関数のローカル変数は、C/C++と全く同じで、スレッド独立だが、
関数引数だけは、スレッド共通だと聞きましたが、これだと再帰コール
できなさそうですが。

**デフォルトの名無しさん** · 2016/03/06(日) 11:18:07.40

あーあまた糞団子が答えちゃった
このスレの事情を知らん奴が見ると余計に迷うぞ

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 11:35:37.89

答えられないクズは存在する価値ないでしょ
とっとと失せろよゴミ

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 11:48:38.53

つーか俺も含めてまだ正解出してないんだけどな

んで正解はこれ、Kepler以降なら再帰に【一応は】対応
http://developer.download.nvidia.com/assets/cuda/files/CUDADownloads/TechBrief_Dynamic_Parallelism_in_CUDA.pdf

誰かツッコミ入れると思ったけど期待外れだったなあ
まあ、俺の説明も何も間違ってないのだけど
(>>862も「出来ない」という説明をしてるのではなく制約を説明しているだけ）

**デフォルトの名無しさん** · 2016/03/06(日) 13:04:30.88

確かにこの団子は鳥屋氏じゃないな。
鳥屋氏はこういう物言いはしない。

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 13:12:58.97

お前は、中村ちゃぷにちゃんが中身の団子が存在すると思ってるんだ
そこから既に間違ってるよ

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 13:18:27.49

ドァンゴが、１３：１８くらいを、お伝えします

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 13:20:55.68

ちゃぷにさんはだんごやさんを嫌ってるはずなので彼がその顔文字を使う理由はないんだよ
自分でNGキーワードに入れてる宣言してたし

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 13:22:13.40

更に言うと彼は精神が壊れて自作PC板で機械翻訳コピペを貼るだけの人に成り下がってしまった

**デフォルトの名無しさん** · 2016/03/06(日) 13:49:39.00

団子＝鳥屋でも別にいいじゃねーか
糞撒き散らすだけならどっちもNGだ

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 13:52:02.23

中村ちゃぷに君と一緒にしないでくれる？

**デフォルトの名無しさん** · 2016/03/06(日) 14:37:46.02

>>879
鳥屋氏に何かあったの? そっちが気になる。

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 14:39:09.93

>>882
だから本人に直接聞いて来いよTwitterアカウントあるだろ

**デフォルトの名無しさん** · 2016/03/06(日) 14:42:36.51

別に本人と直接連絡とれるけどなんか聞きにくいじゃん。
おしえてよ。

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 14:43:25.75

俺なら聞けるのか？ｗ
デタラメ言っちゃうけど？

**デフォルトの名無しさん** · 2016/03/06(日) 14:47:00.65

本人に根掘り葉掘り聞くよりいいよ。

**,,・´∀｀・,,）っ-○○○** · 2016/03/06(日) 14:48:10.37

ちなみに彼は10年以上前から初心者の質問板で「メンヘル屋」と呼ばれてたぞ

**デフォルトの名無しさん** · 2016/03/06(日) 23:01:58.93

日曜日の真昼間に2ch張り付きかよ
さすが糞団子だな

**デフォルトの名無しさん** · 2016/03/07(月) 07:46:57.77

平日の日中ならいいのかｗ

**デフォルトの名無しさん** · 2016/03/07(月) 12:32:22.52

頼むからNGの外に出てくるな>>889＝糞団子

**デフォルトの名無しさん** · 2016/03/07(月) 17:04:45.58

マジで氏ねよ！ここ糞団子！！
お前のせいでスレがめちゃくちゃ！

**デフォルトの名無しさん** · 2016/03/07(月) 19:13:33.67

>>889
なんでそう思ったの？

**デフォルトの名無しさん** · 2016/03/07(月) 21:28:52.80

病気の人が召還してる

**デフォルトの名無しさん** · 2016/03/08(火) 00:17:54.26

糞団子は>>888が効いてる様だなw

**,,・´∀｀・,,）っ-○○○** · 2016/03/08(火) 01:03:04.66

普通ジャン

**デフォルトの名無しさん** · 2016/03/08(火) 06:06:38.08

家でpcにかじり付いてる人なんだろ
病気の人は

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

日曜日に一日中2chが普通なんだ
さすが糞団子

**デフォルトの名無しさん** · 2016/03/08(火) 22:50:34.72

2chやっているおっさんなら起きてから寝るまで2chが普通だろ
いまはスマホで仕事サボって2chしているおっさんが多いからな

**デフォルトの名無しさん** · 2016/03/08(火) 22:59:24.25

>>898
今やちょっとゆる会社だと、トイレが満室になって10分待ちだからな。

**デフォルトの名無しさん** · 2016/03/08(火) 23:40:27.24

そんなにトイレが好きならドアを接着剤で止めてやろうかと思った

**デフォルトの名無しさん** · 2016/03/09(水) 00:03:43.04

平日仕事サボって2chなら、まぁわからなくも無い
理解できないのは日曜日の朝から晩まで2ch張り付きの糞団子

**,,・´∀｀・,,）っ-○○○** · 2016/03/09(水) 00:29:18.66

スマホゲーの体力ゲージ消化してるだけじゃね？
2chなんてとっくにブーム終わってるだろ

**デフォルトの名無しさん** · 2016/03/09(水) 12:01:37.54

>>901
話し相手が欲しいんだよ
スレをウンコ臭くしておいて

**デフォルトの名無しさん** · 2016/03/09(水) 17:46:55.52

へえ、お前、寂しがり屋なんだな

**デフォルトの名無しさん** · 2016/03/09(水) 19:21:18.15

糞団子が急に大人しくなって笑える
かなり効いた様だなw

**デフォルトの名無しさん** · 2016/03/09(水) 19:28:24.12

かまってちゃん
乙

**デフォルトの名無しさん** · 2016/03/09(水) 21:45:15.09

コテ団子をNGにして華麗にスルーしないとね
名無しで煽る時は怒るんじゃなくてファブリーズでも置いておきましょう

**デフォルトの名無しさん** · 2016/03/09(水) 22:07:10.45

そんなにかまってほしいのか

**デフォルトの名無しさん** · 2016/03/09(水) 22:13:39.55

それはお前だろ
さっきからチョロチョロウザい

**デフォルトの名無しさん** · 2016/03/09(水) 22:45:09.81

おまえがそれを言うのか

**デフォルトの名無しさん** · 2016/03/10(木) 00:12:27.45

なんだ
やっぱり、かまってちゃんか？

**デフォルトの名無しさん** · 2016/03/10(木) 05:26:45.85

っ鏡

**デフォルトの名無しさん** · 2016/03/10(木) 20:04:24.43

いや、朝の5時から鏡出されても…

**デフォルトの名無しさん** · 2016/03/10(木) 22:07:55.24

ここはプ板のかまってちゃんが集うところです
Nvidiaユーザーってへんなやつ多いよね。どう？

**,,・´∀｀・,,）っ-○○○** · 2016/03/10(木) 23:24:17.96

ミラーマン！
植草教授！

**デフォルトの名無しさん** · 2016/03/11(金) 00:12:19.52

はい
かまってちゃんの変な奴来ました

**デフォルトの名無しさん** · 2016/03/12(土) 13:14:03.96

誰も居ない？

**デフォルトの名無しさん** · 2016/03/14(月) 00:27:24.09

糞団子が荒らして過疎ったな
ホントろくでもねぇ奴だったな

**デフォルトの名無しさん** · 2016/03/15(火) 19:07:06.06

では素人が質問してもいいですか
オプションのCUDA RuntimeをStatic CUDA runtime library (-cudart static)にしているのにもかかわらずcudart32_75.dllが生成されるんですけど不要ですよね？
消しても動くのでいらなさそうではあるんですけど一応残してます

**デフォルトの名無しさん** · 2016/04/06(水) 05:42:39.04

https://devblogs.nvidia.com/parallelforall/inside-pascal/
https://devblogs.nvidia.com/parallelforall/wp-content/uploads/2016/04/gp100_SM_diagram.png

**デフォルトの名無しさん** · 2016/04/08(金) 23:49:33.98

>>862
この辺りからの流れはとても勉強になる

**デフォルトの名無しさん** · 2016/04/13(水) 00:20:39.50

>>918
CUDAを使ったことがない連中がたくさん来場してたからな
CUDAは映像･画像処理の一般プログラムではなく専門性が非常に高い分野(機械学習や車の自動運転とか)で使うって感じに
なってきたからな。

**,,・´∀｀・,,）っ-○○○** · 2016/04/13(水) 00:37:51.36

レジスタファイルがダイ全体で14Mもあるとかホント化け物ですね
アクセスレイテンシはでかそう

**デフォルトの名無しさん** · 2016/04/13(水) 18:30:01.95

言っているそばからw

**デフォルトの名無しさん** · 2016/04/18(月) 20:53:52.86

以前エクサスケールへむけての改良点てやってたけど
warpサイズが4になったりレジスタファイルの階層化とか
やってるのかな

**デフォルトの名無しさん** · 2016/04/29(金) 07:41:21.25

>>922
機械学習はマルチGPU