【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/

**754** · 2016/03/02(水) 17:56:48.93

あっ

**デフォルトの名無しさん** · 2016/03/02(水) 18:46:16.76

>>775
うーん、速さといってもデータ転送なのかレイテンシなのかで答えようが無いね。
レイテンシに関してはコアに近いところだとCPUの方がクロックが高いから圧倒的にCPUだね。

**デフォルトの名無しさん** · 2016/03/02(水) 19:06:15.16

>>761
いやいや、CUDAユーザーならローカルメモリなんて言い方はしない
そもそもローカルではないし

SRAMはキャッシュとシェアードだけ
ボロ出しまくり糞団子

**デフォルトの名無しさん** · 2016/03/02(水) 20:08:27.14

>>776
なるほど。基本的に残ってないと考えた方がよさそうですね。
バイトニックソートは綺麗に並列化できちゃうので、並び替えは無理っぽいです。
>>778
Occupancy確認しました。ビンゴでした。4%とかなにそれクソ使えてない。
というか、よくよく確認したらsharedもまだまだ余裕があったので、
要素数256でsharedの利用を切り替えていたのを2048まで拡張しました。
8192越えたあたりで悪化し始める、という性質こそ変わりませんでしたが
sharedの恩恵を受けられる範囲が増えたおかげか悪化しつつも
約100万要素でもCPUより7倍強速くすることができました。これでOccupancyは50%。
100%にするにはsharedのサイズの設定を変えなきゃならないけどやり方わかんない…。
しかし、こんなアホみたいな理由で遅くなっていたなんて。
みなさまにはご迷惑おかけしました……。
>>780
やっぱりCPUですか。GPUの性能をフルに発揮させるには
それなりに大きい規模を並列化して誤魔化さないといけないわけですね。

**754** · 2016/03/02(水) 20:12:34.57

名前入れ忘れたりsage忘れたりしてますが>>782は754です。
IDで察してくださいすみません。

**デフォルトの名無しさん** · 2016/03/02(水) 20:36:32.21

迷アドヴァイスを見るに見兼ねたスレ住民のおかげで助かりましたとさめでたしめでたし

**デフォルトの名無しさん** · 2016/03/02(水) 21:24:19.27

>>782
シェアードをローカルとか言う奴の事は聞かないように
GPUはシェアードメモリを使ってなんぼだから

**デフォルトの名無しさん** · 2016/03/02(水) 21:47:03.61

っていうかgpgpuでソートが
そもそも間違ってるよなあ

**,,・´∀｀・,,）っ-○○○** · 2016/03/02(水) 23:53:33.06

てかFermi以降のSMはキャッシュを構成するSRAMの一部を
GMとは独立した空間のスクラッチパッドメモリとして割り当ててるだけなんだけどね
SMXが実質的なコアとみなすならシェアードメモリって別にシェアードでもなんでもないぞ

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

ますます支離滅裂だな
シェアードメモリがシェアードじゃないとかCUDAの概念を全否定する発言じゃねぇか
つまりお前はCUDAの事なんか全く判ってない
ここはCUDAのスレだぞ
そろそろ出て行けや糞団子

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 00:24:01.33

まったく技術論になってないぞ
NVIDIAの用語がインチキ用語なのは今に始まったことじゃない
AMDのBulldozerのコア詐称以前からのコア数水増しやってるのがNVIDIA

命令ストリームレベルで見ればSMXが本来のコアで
CUDA Coreと呼んでるものはSIMDユニットの1エレメントにすぎない

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

確かにCUDAというかNVIDIAの定義しているシェアードメモリは一般的なシェアードメモリとは違うね。どちらかというとローカルメモリといった方が正しいかもしれん。

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 01:01:39.27

同じコアだからこそメモリを介さずレジスタ上でシフト・シャッフルができる（Kepler以降）

CUDAのSIMTとは単純に32要素1024ビットのSIMDを32スレッドと読み替えてるだけ
この点はG8x世代から変わらないね

**デフォルトの名無しさん** · 2016/03/03(木) 01:48:33.97

自演会話すれば糞団子が味噌団子に変化するじゃないぞ？団子さんよお

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 07:58:21.23

自演認定しても頭がよくなりませんよ、他所者君
CUDAをちゃんと勉強してる人にはみんな知ってる事実だ

**デフォルトの名無しさん** · 2016/03/03(木) 08:43:19.76

また夜中の自演が始まったか
バレてないと思うところが糞団子だね
何度も言うがここはCUDAのスレ
僕の考えたシェアードメモリ論ならチラシの裏に書いとけ

**デフォルトの名無しさん** · 2016/03/03(木) 09:02:21.96

CUDAでいうシェアードメモリはCPUでいう一次キャッシュだし、CELLで言うところのLSだから、SMメモリとかコアメモリとかにしておけばよかったのにね。
CUDAの黎明期はGPUをスパコンって呼ぶ戦略があったから、1スレッドを一コアに見立て、共有メモリだと言っていたのかもしれないね。

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

>>795
それを言うならオンチップメモリ
少なくともソフトレイヤーではキャッシュとシェアードは別物
そもそもMaxwellではデータバスも異なる
んでここはCUDAのスレだ
スペオタと糞団子は去れ

>>792
cudaDeviceSetCacheConfigは試してみたのかな？

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

>>792
俺に聞かれてもなｗ
Occupancy確認してからスムースになってきたみたいだし、使わなくても順調ならそれでいいという
考えもある

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 12:34:58.87

ソフトレベルって言ってもなぁ
.cuやPTXのレベルならそうだろうがcubinレベルでみて言ってるんじゃないんだろ？

実際Sharedなんて言ってもSMX（実質上のコア）の中でしか参照できないんじゃ名前に偽りありとしか
SPMDをオレオレ用語で置き換えた概念で説明すればSharedなんだろうが物理的にはCellのSPUのLSと同じトポロジにすぎん

**デフォルトの名無しさん** · 2016/03/03(木) 14:52:54.35

逆に複数のsmx間で共有されてたら
くそ遅いだろう
ローカルだから他のsmx気にしなくていいわけで

**デフォルトの名無しさん** · 2016/03/03(木) 14:54:11.01

>>798
相変わらず頓珍漢なこと言ってるな
お前がどう思うかは全く興味がないから
ここはCUDAスレだ
CUDA使った事ない糞団子は去れ

**754** · 2016/03/03(木) 14:57:08.65

>>796
CUDA_Occupancy_Calculator.xlsで確認してたんですが、
CC3.0だと48kb/16kbと16kb/48kbの他に32kb/32kbがあるっぽくて、
32kbなら100%になる、みたいな謎結果が出てたんです。
で、うちの古い参考書にはその設定のやり方が載ってなかったんです。
だけどそのあと.xlsファイルを開き直したらなぜが100%になってました。
なんでやねん。

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 15:08:31.08

>>800
お前がCUDA使ったことないゴミだと自己紹介してることはわかったよ
俺は初代の8系から触ってるんだがな
墓穴ほり過ぎ

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

また低脳の団子が自演かよ!?
いい加減にして欲しい
コイツがいるとスレが腐るんだけど（怒）

**デフォルトの名無しさん** · 2016/03/03(木) 19:52:45.84

>>802
初代から使ってそのレベルならお前には向いてないからもう諦めろ
結局質問者には何ひとつ有益な情報を与えられない低レベル糞団子

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

>>801 （>>754）
ああいう大人（団子）みたいになったらだめですよ（笑い

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

団子ちゃんは詳しいと思う。

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

文句言うだけの奴より団子のほうがこのスレの役に立っている。

**デフォルトの名無しさん** · 2016/03/03(木) 21:01:42.18

　　　　　　　　　　　　　　　【親米涙目】　★　　　ロシアがアメリカを一発ＫＯ　　　★　【反中逃亡】

　　宇宙人側からの申し入れは、とにかく核の利用と戦争をやめなさい、もう一つは宇宙人の存在を公表しなさい。
　つけ足したのが、６０年の猶予を与えましよう、２０１４年には発表しなければならなかったんですね、宇宙人の存在を。
　　　　　ロシアという大国の首相がね、あれは冗談だよでは済まないですね、しかも２回も言ってるんだからね。
　　　　　　　　　　　　　　　　　　　　　https://www.youtube.com/watch?v=FIRXKetUkq8

　　　　　　　　　　　【親米屈辱】　★　　　宇宙人いない派のおバカさん残念ｗｗｗ　　　★　【反中赤面】

　　　　　　NASAは、UFOに対して長年取ってきた態度のために、無用な組織とされることを恐れています。
　　　　　　　　　　　　　　　　　マイトレーヤが公に現れるにつれて、UFOが姿を表すでしょう。
世界平和の脅威は、イスラエル、イラン、アメリカです。イスラエルの役割は跪いて、パレスチナに許しを請うことです。
　　　　　　　　　　　　　　　　アメリカによる他国の虐待に反対の声を上げなければなりません。
　　　　彼らは今世紀（21世紀）をこの帝国が出来上がるアメリカの世紀と呼ぶ。しかし、そうはならないだろう。
　　彼らが世界中に‘民主的’制度を確立したいという衝動をコントロールするのは、マイトレーヤの任務である。
Q 経済崩壊が2015年から、テレパシー世界演説が2017年の初めなら、2016年に大戦でしょうか。A 大戦は起こりません。

　　　【親米命乞い】　★　　　マイトレーヤのテレパシーによる世界演説は英国ＢＢＣが放送　　　★　【反中土下座】
　　　　　　　　　　　　　　　　　　　　　https://www.youtube.com/watch?v=6cOvo6n7NOk

　　　　　　　　　　　　　　　　マイトレーヤが世界に向かって話をする準備は良好に進行している。
　　25分か35分くらいかもしれませんが、歴史上で初めて、世界的規模のテレパシーによる接触が起こるのです。
　　テレビ中継はこの出来事のために存在するのであり、この時、初めてご自身の本当の身分を明らかにされます。

**デフォルトの名無しさん** · 2016/03/03(木) 21:02:13.41

はい、単発来ましたw

糞団子のどこが有益なんだ？
シェアードメモリの割り当て変更にも答えられずに、ローカルメモリとか嘘の情報垂れ流してる奴だぞ？
まともなCUDAユーザーならローカルメモリとシェアードメモリが別物なんて事は常識以前の問題だ
糞団子の自論なんか誰も聞きたくない
嫌われ者の糞団子は去れ！

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 22:03:27.41

他人のレスで我が物顔してるだけじゃん
無益な厨房だな

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 22:12:39.51

つーか自称シェアードメモリが実質的なLocalStorageでしかないことは知っておかないとハマる事例もあるんだけどね
SIMTはAVXユニットを8コア、Xeon Phiの1コアを16コアって読み替えるトンデモ抽象化

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

自演するな！糞団子

**,,・´∀｀・,,）っ-○○○** · 2016/03/03(木) 22:34:50.71

別に何も間違ったことは書いてないんだが理解できないバカがいる

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

シェアードメモリをローカルメモリとか言う嘘を垂れ流してる事はスルーなんだな
簡単な初期化関数すらアドバイス出来ない低レベル糞団子

ちなみにシェアードメモリは自称でもなんでもなく公式な名称だ
気に入らないなら自分で会社起こして俺の考えた最強のGPUでも作ればいいさ
まぁ、お前のレベルじゃディスクリートすら無理な事は明らかだがな
早く去れよ糞団子

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

さっき、Anacondaの最新バージョンである2.5.0(64-bit　python3.5.1)を入れたの
だけど、それにコマンドプロンプトがデフォルトで入っていないので、pip installが
出来なくて困っています。ipythonやインタラクティヴシェルは入っています。

前のAnacondaのバージョン(python2.7で使っている)であれば、ipythonや
対話シェルからq()やquit()で、普通に、コマンドプロンプトに戻れて、そこから
pip installなどが出来たのになぜなのでしょうか。

普段使っているコマンドプロンプトはpython2.7用になっているので、そこから
python3.5.1用のpip installも出来ません。コマンドプロンプトをpythonのヴァージョン
ごとに切り替えることができるのかもしれないけど、その方法も分かりません。

どなたか、知っている人がいたら教えて下さい。
よろしくお願いします。m(__)m

**デフォルトの名無しさん** · 2016/03/03(木) 23:25:12.80

>>815
恐ろしくスレ違いなんだけど
python2系と3系を共存させるならパスを適切に設定しないとダメだよ
詳しくはググってくれ

取り敢えずコマンドプロンプトでpyって打ってからpipしてみたら？

**デフォルトの名無しさん** · 2016/03/03(木) 23:27:56.45

>>809
どう違うの？

**デフォルトの名無しさん** · 2016/03/03(木) 23:45:09.31

>>817
なぜググらない？
http://www.gdep.jp/page/view/252

**デフォルトの名無しさん** · 2016/03/04(金) 00:05:36.72

ha?

**デフォルトの名無しさん** · 2016/03/04(金) 00:13:03.67

>>816
パスは、最新のAnacondaをインストールする時に、パスを自動設定するなら
チェックを入れよという欄があって、チェックを入れてインストールを終了
すると、新しい環境設定が入っています。

念のためシステムで環境設定でパスをチェックしたけど、Anaconda3で、
新しいのが入っている。Anacondaは、3ヴァージョン使っている
のだけど、新しくなるほど使いにくくなっている感じですね。

q()でコマンドプロンプトに戻れると>>815で書いたのは、間違えだった。
それは、R言語のやり方だった。

**デフォルトの名無しさん** · 2016/03/04(金) 00:16:41.20

nv用語だとローカルにないものが
ローカルメモリなんだなこりゃややこしい

**デフォルトの名無しさん** · 2016/03/04(金) 00:25:56.51

今更の問題だけど、pythonはやはり、2系と3系の互換・移植性が
悪いのが問題かな。今日買った、

基礎 Python (IMPRESS KISO SERIES) 単行本（ソフトカバー） – 2016/3/4

大津真 (著)

ttp://www.amazon.co.jp/%E5%9F%BA%E7%A4%8E-Python-%E5%A4%A7%E6%B4%A5-%E7%9C%9F/dp/484438015X

が、3系で、ちょっと変わった雰囲気のコードの使い方があったので、面白いの
だけど、2系と併用だと>>815みたいになって、pip　installも出来ないし、
ipythonでしか、3系が使えないのが困りもの。2系で使える3系のものも
あるけれど…。

2系を全部消せば3系を使えるのは分かっているのだけど、2系の方が
出来ることがまだまだ多い感じがするからね。まだ2系が多数派でしょ。

あ、仮想環境ヴァージョンを切り替えてやればいいのかな。

**デフォルトの名無しさん** · 2016/03/04(金) 00:38:24.47

あ、pythonのクダスレだと、勘違いしてました。
失礼いたしました。

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

>>821
何を今更
だからこそ嘘を垂れ流す糞団子はここから去れよ

**,,・´∀｀・,,）っ-○○○** · 2016/03/04(金) 01:07:12.95

ウソ扱いしてるお前が自覚のないウソつきなんだけど

**,,・´∀｀・,,）っ-○○○** · 2016/03/04(金) 01:15:03.88

>>821
もっとも、スピルされたレジスタのデータが真っ先に格納されるのはGDDR*ではなく
キャッシュだけどな
なぜならキャッシュはDRAMの読み書きを高速化するためにあるから

**デフォルトの名無しさん** · 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板で機械翻訳コピペを貼るだけの人に成り下がってしまった