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

**デフォルトの名無しさん** · 2016/04/29(金) 19:49:33.11

Pascalは忘れられてるけど、倍精度も強化されてるんだよね
単精度を徹底的に削った思い切った方針には目を見張った

**,,・´∀｀・,,）っ-○○○** · 2016/04/29(金) 21:17:01.56

大型スパコン案件に特化して設計したものをGPUに流用してるだけで
それが全部だと言い張られても困りますが

**デフォルトの名無しさん** · 2016/04/29(金) 22:49:34.54

なんでおまえが困るんだよw

**デフォルトの名無しさん** · 2016/04/30(土) 00:41:27.46

ＧＰＧＰＵは深層学習のブレークで今後一気に需要が高まる？

**デフォルトの名無しさん** · 2016/04/30(土) 11:52:36.46

高まるだろうね
倍精度と半精度がちょうどいい具合

そうやってGPGPU市場を頑張ってもらって12.1の妨害行為は止めて欲しいわ

**デフォルトの名無しさん** · 2016/04/30(土) 22:20:59.01

その需要の大半はnVIDIAが発表したEyerissみたいな専用チップに喰われるのでは

**,,・´∀｀・,,）っ-○○○** · 2016/04/30(土) 22:33:42.88

あれはCNNに特化したものでGPUほどのプログラマビリティはない
ディープラーニングって一言で言っても1個のASICで賄えるほど狭い分野じゃない

**デフォルトの名無しさん** · 2016/04/30(土) 22:45:19.47

CNNに特化て具体的には行列演算てこと？

**,,・´∀｀・,,）っ-○○○** · 2016/04/30(土) 23:20:42.71

疎行列に特化したマトリクス演算回路とそれなりの帯域のリングバスネットワークが主な構成要素
って覚えておけば大体間違いない
実はGPUの特性と正反対（GPUは密行列なら得意）

**デフォルトの名無しさん** · 2016/04/30(土) 23:25:48.76

fermi世代ってdouble⇒floatで倍速なはずだよね？
ちっとも速くならないんだが。

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

大分前だけど構造体でデータを渡すとdouble扱いになったことはある
原因は不明

**デフォルトの名無しさん** · 2016/05/01(日) 10:04:45.15

>>937
thx。

調べてみた感じだと、倍精度が無いカードで暗黙のdouble→floatキャストがあったり、
floatへの数値代入で接尾語Fを付けないとdoubleにキャストされたりもあるっぽいね。

fermiでarch=compute_20,sm_20指定してるから基本的には宣言通りだと思ったのだが、実際どうなってるのか・・。
全部floatに置き換えてもわずかしかスピードアップしないのは、バイトコピー分が浮いてるだけな気がする。

**デフォルトの名無しさん** · 2016/05/01(日) 10:24:43.67

匿名通信（Tor、i2p等）ができるファイル共有ソフトBitComet（ビットコメット）みたいな、
BitTorrentがオープンソースで開発されています

言語は何でも大丈夫だそうなので、P2P書きたい！って人居ませんか？

Covenantの作者（Lyrise）がそういう人と話したいそうなので、よろしければツイートお願いします
https://twitter.com/Lyrise_al

ちなみにオイラはCovenantの完成が待ち遠しいプログラミングできないアスペルガーｗ

The Covenant Project
概要

Covenantは、純粋P2Pのファイル共有ソフトです

目的

インターネットにおける権力による抑圧を排除することが最終的な目標です。そのためにCovenantでは、中央に依存しない、高効率で検索能力の高いファイル共有の機能をユーザーに提供します

特徴

Covenant = Bittorrent + Abstract Network + DHT + (Search = WoT + PoW)

接続は抽象化されているので、I2P, Tor, TCP, Proxy, その他を利用可能です
DHTにはKademlia + コネクションプールを使用します
UPnPによってポートを解放することができますが、Port0でも利用可能です(接続数は少なくなります)
検索リクエスト、アップロード、ダウンロードなどのすべての通信はDHT的に分散され、特定のサーバーに依存しません
ｈ

**デフォルトの名無しさん** · 2016/05/01(日) 11:36:33.54

eyerissは大量のデーターが複数のフィルター処理されるのに向いてる

**デフォルトの名無しさん** · 2016/05/09(月) 19:46:40.50

質問があります
環境はwindows 10 | arch linux, cuda 7.5, gcc | visual c++ 2013です

#include <stdio.h>

__global__ void a(char *A){
}

int main(){
a<<1, 1>>(); // (1)
}

(1)で「expression must have integral or enum type」とエラーが出て困っているのですが
ホストコンパイラを呼び出す際に何か間違えているのでしょうか

コンパイルに使用したコマンドは
windows
　nvcc -ccbin "C:\Program Files (x86)\Microsoft Visual Studio 12.0\VC\bin" -c -DWIN32 -D_CONSOLE -D_MBCS -Xcompiler /EHsc,/W3,/nologo,/Wp64,/O2,/Zi,/MT test.cu
linux
　（失念したが大体同じ）
です

**デフォルトの名無しさん** · 2016/05/09(月) 23:47:16.41

>>941
≪≫のところじゃない？
二重カッコじゃなくて、三重カッコにすべきかな

二重カッコだと別の演算子になって、意味が変わってくる

**デフォルトの名無しさん** · 2016/05/09(月) 23:54:10.06

ああ、3重括弧でしたか
間抜けなミスで騒いでしまって申し訳ないです
ありがとうございます

**デフォルトの名無しさん** · 2016/05/20(金) 21:20:51.03

グーグルがTensolFlow用のTPUとかいうの作ってたらしいがどうなの？

**sage** · 2016/05/23(月) 17:33:55.45

CUDA fortranで一つの配列に対して2GB以上のメモリを
アロケートすることが出来ず、
PGIのサイト見たら「2GBの壁」とかって説明されているのですが、
これってどうにか克服出来ないですかね？

**デフォルトの名無しさん** · 2016/05/23(月) 18:16:51.67

>>945
何gbくらい必要なの？

**デフォルトの名無しさん** · 2016/05/23(月) 18:47:01.54

>>945
fortranなんか使っているからだよ。

**デフォルトの名無しさん** · 2016/05/23(月) 19:56:12.00

３２ビットの壁ってこと？

**デフォルトの名無しさん** · 2016/05/23(月) 22:17:34.30

CU_LIMIT_MALLOC_HEAP_SIZE か cudaLimitMallocHeapSize を拡げりゃいいんじゃないかな。

**デフォルトの名無しさん** · 2016/06/01(水) 19:25:08.39

ttps://developer.nvidia.com/cudnn
>Supported on Windows, Linux and MacOS systems with Pascal, Kepler, Maxwell, Tegra K1 or Tegra X1 GPUs.
のTegra K1の文字を見て、サポート復活を喜んでいたのですが、
Tegra K1は32bitでして、
必要なCUDA7.5は64bit必須
もしかして、cuDNNとCUDAは無関係で独立して使えたりしますか？

**デフォルトの名無しさん** · 2016/06/03(金) 02:14:25.16

CUDA APIで１番始めに呼んだAPIは時間がかかるとのことですが、
時間計測に影響が出ないように事前にどんな関数を呼んでますか？
慣習だったり、よく使われる関数があれば教えてください。

**デフォルトの名無しさん** · 2016/06/04(土) 23:09:31.90

工学じゃなくて物理系でもcuda覚える価値ってある？

**デフォルトの名無しさん** · 2016/06/04(土) 23:54:08.10

安く実験したいなら
N体や粒子法なんかは普通のCPUじゃたいしたことはできない

**デフォルトの名無しさん** · 2016/06/05(日) 01:52:52.19

むしろ、科学計算したい人のための技術だと思う

**デフォルトの名無しさん** · 2016/06/12(日) 02:11:09.53

他人にも聞こえんの？

**デフォルトの名無しさん** · 2016/06/12(日) 12:50:48.33

この部分がエラー吐くんだが
cuda7.5

struct cuComplex {
float r;
float i;
cuComplex( float a, float b ) : r(a), i(b) {}
__device__ float magnitude2( void ) {
return r * r + i * i;
}
__device__ cuComplex operator*(const cuComplex& a) {
return cuComplex(r*a.r - i*a.i, i*a.r + r*a.i);
}
__device__ cuComplex operator+(const cuComplex& a) {
return cuComplex(r+a.r, i+a.i);
}
};

**デフォルトの名無しさん** · 2016/06/12(日) 13:24:56.47

CUDA by EXAMPLEのサンプル？
<cuComplex.h>の定義とぶつかってるんじゃね？

**デフォルトの名無しさん** · 2016/06/12(日) 13:39:51.78

>>957
そうそう
エラーはこんな感じ

Documents/Visual Studio 2013/Projects/template/cuda_by_example/chapter04/julia_gpu.cu(43): error
: calling a __host__ function("cu_Complex::cu_Complex") from a __device__ function("julia") is not allowed

**デフォルトの名無しさん** · 2016/06/12(日) 13:50:41.56

試してないけど、コンストラクタに__device__付けてないのが悪いと言っているような。

**デフォルトの名無しさん** · 2016/06/12(日) 14:09:28.62

>>959
ありがとうございます
解決しました

**デフォルトの名無しさん** · 2016/06/15(水) 14:42:46.55

【GPGPU】くだすれCUDAスレ part8【NVIDIA】
http://echo.2ch.net/test/read.cgi/tech/1416492886/

**デフォルトの名無しさん** · 2016/06/15(水) 14:44:14.63

誤爆

【GPGPU】くだすれCUDAスレ part8【NVIDIA】
http://echo.2ch.net/test/read.cgi/tech/1465969275/

**デフォルトの名無しさん** · 2016/06/16(木) 14:11:42.71

615 ：名前は開発中のものです。：02/06/10 21:48 ID:???
>将来的にはスカートの頂点を触ってヒラヒラさせたりしたいので、どうせシステムメモリーに
ボーンでやる気なのか？つらそー。

616 ：◆CKmHzWpc：02/06/11 03:48 ID:zXaLQtp.
みなさんありがとうございます
１．モデルを分割するのが嫌→頂点シェーダ－を使うorシステムメモリに置いてソフトでやる
２．モデルを分割してもいい→ハードのジオメトリブレンドを使ってやる・・
というのが考えられる・・・んですよね
色々組み合わせを試して、高速なのを選びたいと思います

>>615
これは、スカートの頂点をCPUで直接操作しようと思っています
ボーンでやらずに、簡単なクロスシュミレーションのような事を目指しています
どうせモデルの一部をシステムメモリに置くなら、もう全部システムメモリにし
てもいいか、と思ったのです

**デフォルトの名無しさん** · 2016/07/20(水) 14:03:52.95

tess

**デフォルトの名無しさん** · 2016/09/02(金) 13:17:38.94

erlangならね