【GPGPU】NVIDIA CUDA質問スレッド

【GPGPU】NVIDIA CUDA ..

892:デフォルトの名無しさん
08/12/18 00:49:51
URLﾘﾝｸ(sourceforge.net)
どうぞ

893:デフォルトの名無しさん
08/12/18 01:47:35
>>892
それ使ってるんですが・・・readmeを嫁ということですね

894:デフォルトの名無しさん
08/12/18 03:39:45
ptx出力がどうなるように最適化していけばいいんですか？

895:デフォルトの名無しさん
08/12/18 07:42:38
>>894
>889なら、registerの数を減らしたいんだろ。ptxの出力を読んで少なくなっていればOKだw

896:デフォルトの名無しさん
08/12/20 00:21:41
>>862
NVIDIA、モバイルGPUドライバを配布開始
～ノートでPhysXなどCUDAアプリケーションが利用可能に

12月18日(現地時間)配布開始

　米NVIDIAは18日(現地時間)、モバイルGPU用の汎用ドライバを同社ホームページ上で配布開始した。
対応GPUはGeForce 8M/9M、およびQuadro NVS 100M/300Mシリーズ。

　従来、NVIDIA製モバイルGPUのドライバは、ノートPC本体のメーカーが提供していたが、
今回よりNVIDIAが汎用ドライバを提供することになった。対応OSはWindows XP/Vista。

　バージョンはデスクトップ版の180代に近い179.28 BETAで、CUDAに対応。これにより、PhysXや、
CUDA対応の動画編集ソフト、Photoshop CS4などでアクセラレーションが効くようになる。
ただし、PhysXについてはビデオメモリが256MB以上必要となる。

　なお、ソニーVAIO、レノボThinkPad、デルVostro/Latitudeシリーズ、およびHybrid SLI構成の
製品については対象外となっている。

　現バージョンはベータ版で、WHQL準拠の正式版は2009年初頭の公開を予定している。

URLﾘﾝｸ(pc.watch.impress.co.jp)

897:デフォルトの名無しさん
08/12/21 00:11:21
なんでvaioは対象外ですか？いじめですか？？

898:デフォルトの名無しさん
08/12/21 07:03:38
vaioとか一部のノートはリコールされてるから

899:デフォルトの名無しさん
08/12/21 19:00:56
それ以前にノートだとメーカーが色々やってるから
一般的にドライバ更新さえノートのメーカー対応待ちだろ。

900:デフォルトの名無しさん
08/12/21 22:30:31
シェーダプロセッサってエラーがあってもゲームでは無視されて
認識出来ないレベルで画面が崩れるだけだから
一見正常に動いてると思っててもＣＵＤＡは動かない人が多いかも
ドライバが対応してるとか以前にハードが用件を満たしてない

901:デフォルトの名無しさん
08/12/24 10:44:05
どこでエラーになっているのかデバッグしろよ

902:デフォルトの名無しさん
08/12/25 15:53:55
ちょいちょいCUDAをいじり始めてるんですが、OpenCLの規格が決まったことで、
CUDAの仕様も大幅に変わってしまうんですかね。

903:デフォルトの名無しさん
08/12/25 18:15:24
MacBookProなんかだと自前でドライバ突っ込めばCUDA使えたんだけど，
VaioやThinkPadだとどうですか？
使っている人がいれば情報求む．

904:デフォルトの名無しさん
08/12/25 18:21:26
5秒制限の条件とか回避の仕方とかがよくわからないのですが
教えてもらえないでしょうか？

905:デフォルトの名無しさん
08/12/26 01:44:56
ThinkPad T61はOpenGL関係以外は使えた。

906:デフォルトの名無しさん
08/12/26 03:48:04
OpenCLはAPIのみだからCUDAとは別にライブラリが用意されるだけだと思うよ
こっちは専用コンパイラも必要ない

907:デフォルトの名無しさん
08/12/26 14:22:18
>>904
Gridを何回も起動するしかないのでは。だから、途中結果をメモリにのこしとかないといけない。
面倒だよね。

908:デフォルトの名無しさん
08/12/26 14:35:42
>>904
Xを使わない、グラフィック機能を使わなければいけるはず。

…何のためのグラフィックカードやねんって感じだが

909:デフォルトの名無しさん
08/12/26 17:45:27
>>907 >>908
やっぱりそれぐらいしかないですか。
現状でグラフィック機能止めて使うのは無理なので
Gridの分割を考えてみます。

返答ありがとうございます。

910:デフォルトの名無しさん
08/12/26 18:15:35
Vistaだとレジストリ弄ればもうちょっと長くできるよーみたいなこと言ってたような、言ってなかったような

911:デフォルトの名無しさん
08/12/26 22:25:09
使うグリッド数を制限すればなんとかならなくもないんだが、効率は落ちるんだよね。
つーか、効率重視ならもう一枚挿すしかない希ガス。

912:,,・´∀｀・,,）っ-○◎●
08/12/27 00:47:07
こんなの教えて貰った
URLﾘﾝｸ(compview.titech.ac.jp)

913:デフォルトの名無しさん
08/12/27 05:05:39
どのみちCUDA実行中は画面の描画が止まるから割り込み入れる工夫は必要でしょ

914:デフォルトの名無しさん
08/12/27 08:47:20
描画要求の方がGridよりも優先されているようにしか見えないんだけど。
止まってくれたらどんだけ効率が改善することかw

915:デフォルトの名無しさん
08/12/27 11:52:08
>>912
面白そうだけど、東工大は遠いなぁ。ワープアSEに新幹線は高い。
わがまま言うと、神戸でやって欲しい。

916:デフォルトの名無しさん
08/12/27 14:13:48
大岡山の東工大なら歩いて行けるぜ！

917:,,・´∀｀・,,）っ-○◎●
08/12/27 23:27:30
>>915
六甲台か？
あの急勾配の坂は登りたくないな。

918:デフォルトの名無しさん
08/12/29 05:06:23
>>912
これは紹介レベルだから、たいした内容ではないよ

919:デフォルトの名無しさん
08/12/29 11:44:34
CUDAをクアッドコアで動くように出来ればいいのに

920:デフォルトの名無しさん
08/12/29 12:15:26
>>919
どういう意味?

CUDAをQuadCoreCPUで動かすことなら何の問題もなくできるのだけど。

921:デフォルトの名無しさん
08/12/29 12:44:32
っOpenCL

922:デフォルトの名無しさん
09/01/05 20:09:11
VIPPERの諸君　立ち上がれ！
以前韓国を打ち破ったのは誰か知らない、小学生や、他サイトの管理人がVIPを馬鹿にしている！

全員集結せよ。そして全員で打ち倒すのだ！

もし時間がない人は、この文章をそのまま他のスレに貼ってほしい！

対策本部
ｽﾚﾘﾝｸ(news4vip板)

今こそ　Vipperの意地を見せつけるのだ！

923:デフォルトの名無しさん
09/01/08 17:49:55
nc = 100*100
bs = 50
dim3 dimBlock(bs,bs)
dim3 dimGrid(sqrt(nc)/dimBlock.x,sqrt(nc)/dimBlock.y)
kernel<<<dimGrid, dimBlock>>>(idata, odata, sqrt(nc)

__global__ void kernel(float* idata, float*, odata, int nc)
{
index=blockIdx.x * blockDim.x + threadIdx.x +
+(blockIdx.y * blockDim.y + threadIdx.y) * nc
}
この時のイメージは、Gird:2x2、Block:50x50でよいのでしょうか?
それとこのままグローバルメモリで計算するのはできるのですが、一旦
シェアードメモリに退避して計算してグローバルメモリに戻す方法が
サンプルを見てもうまくいきません。どういう感じになるのでしょうか?

924:デフォルトの名無しさん
09/01/08 18:11:37
いつも疑問に思いつつ使っているのですがグローバルメモリ使うときも
共有メモリ使うときも
float* data;
CUDA_SAFE_CALL(cudaMalloc((void**)&data, sizeof(float) * 100)
CUDA_SAFE_CALL(cudaMemcpy(data1, h_data, sizeof(float) * 100, cudaMemcpyHostToDevice))
になるのでしょうか?これはグローバルメモリにとっているのですよね。
共有メモリの話は、kernelのほうで使用するかどうかですよね?

コンスタントメモリ使うときだけどうも
CUDA_SAFE_CALL(cudaMemcpyToSymbol(data, h_data, sizeof(float) * 100))
としてホストからコピーするんですかね?

925:デフォルトの名無しさん
09/01/08 23:27:14
>>924
cudaMemcpyToSymbolは定数メモリの他、グローバルメモリでも使えるよ。
要は、device空間で定義されたメモリと言うニュアンスでしかない。
逆に言えば、cudaMemcpyToSymbolを使わないとhost側のメモリ空間に割り付けられたアドレスを利用しようとしてしまう。
# この辺が判り難いのよね。

共有メモリはグリッド実行中しか使えないからhost側からは為す術がない。

>>923
どのサンプル見てどうやったらどう巧くいかないのか具体的に。
少なくとも、>923のような処理なら共有メモリの使いようがないんだけど。

926:デフォルトの名無しさん
09/01/09 01:50:13
warp単位32スレッドで同じ命令を行って、
実際に一度に動くのは8スレッド

こういう考え方でいいのですか？

927:デフォルトの名無しさん
09/01/09 18:21:54
>>925
結局現状では、cudaMemcpyXXXの違いは理解できそうにないです。
まあそれはよいとして(全然良くないですが)
>>923
のidataは、実は、一次元配列なんです。なのにカーネルには2次元で利用しよう
としているんです。
float host_idata[nc]
float* idata
cudaMalloc((void**) &idata, sizeof(float) * nc)
cudaMemcpy(idata, host_idata, sizeof(float) * nc, cudaMemcpyHostToDevice)
こんな感じのグローバル変数がidata。
としたとき、ncとbsの設定値によってはうまく処理できないことがあるのです。
これが謎です。わかる方いませんか?もちろんncはsqrtで整数になる値、設定する
値は、割り切れる値です。nc=1000*1000,bs=500で処理結果がおかしくなっています。
kernelの読んだあとエラーでinvalid configuration argument.になってます。

それと、共有メモリのほうですがこれは分かりました。
kernel<<<dimGrid, dimBlock, sizeof(float) * nc>>>(idata, odata, sqrt(nc))
第3引数の共有メモリのサイズを定義しないとだめだった。ちゃんと説明書を
読んでいなかったというレベルでした。これは解決です。

928:デフォルトの名無しさん
09/01/09 19:05:25
>>927
ブロックは512x512x64と出ているけどいちお3次元は分からないので512x512
としてxで512使い切るとどうもyではもう1しかとれないようです。
ncは、そのままだとエラーが出ないレベルは、bs=20つまり20x20=400ブロック
使用。21,22は割り切れないので設定できない。みたいです。
とすると一次元配列を使用するのにわざわざ、x,yの次元のブロックを使いきる
テクニックは必要なさそうですね。と推察できると思います。結果からですが
ただ、これがバグであるなら次のCUDAのバージョンアップで使えるようになる
かもですね。

929:デフォルトの名無しさん
09/01/09 20:53:02
>>928
次元とは別に1ブロックあたり512スレッドが上限じゃなかったっけ？
どっかにきちんと書いてあったはずだからバグではないよ。

930:デフォルトの名無しさん
09/01/09 22:23:41
私は1次元でやってしまうことが多いなぁ。
で、スレッド数についてはdeviceQuery参照で。

931:デフォルトの名無しさん
09/01/10 21:32:30
スレッドの最大数はAPIで取得できる。全部512だと思う。
ブロックの最大は65535だから65535*512以上は1次元のグリッドでは扱えない。
あと最大は512だけど、192か256くらいのほうがパフォーマンスが出る。

932:デフォルトの名無しさん
09/01/14 01:57:36
CUDA 2.1 Release
URLﾘﾝｸ(forums.nvidia.com)

933:デフォルトの名無しさん
09/01/14 21:34:37
情報サンクス。早速、Linux版ゲットした。

934:デフォルトの名無しさん
09/01/15 01:08:35
9600GTでVideoDecodeのサンプル動かそうとしたら正常に動かなくて、
結局SDK2.0+178.24まで戻さないとダメだった。

2.1にもソースは付いてるけど、2.0のままみたいだね。

935:,,・´∀｀・,,）っ-●◎○
09/01/15 07:55:42
>>931
1 warp = 32 threadで、GeForce 8/9が24warp/block、GT200で32warp/blockが最大だから
768thread/blockと1024thread/blockが最大なんだけどね。本来は。
CUDAのドライバ側が512でリミッタかけてるんだ。罠としか言いようが無い。
逆に言うとCUDAを経由しなきゃ目いっぱい使えるかもね。

ただ、スレッドインターリーブすると１スレッドあたりで使えるレジスタ本数が減っちゃうんだよね。
メモリレイテンシを隠蔽するならスレッドを目いっぱい使ったほうがいいし
逆に一時変数を何度も再利用する場合は、thread/blockを減らして１スレッドあたりの仕事量を増やしたほうがいい。

936:デフォルトの名無しさん
09/01/15 15:42:19
CUDAってシーユーディーエーって読んでいいの？
それともクダとか、なんか読み方あったりする？

937:デフォルトの名無しさん
09/01/15 15:57:04
>>936
クーダ

938:,,・´∀｀・,,）っ-●◎○
09/01/15 15:58:34
ネイティブの人は「クーダ」って発音してたよ。
Cellは「くた」だから注意な。

939:デフォルトの名無しさん
09/01/15 16:01:48
ありがとう！明日から会社でクゥゥゥウーダ！！って叫びまくるわ

940:デフォルトの名無しさん
09/01/15 16:02:44
きゅーだ　って発音してる俺は異端ですかね？

941:デフォルトの名無しさん
09/01/15 16:23:46
>940
そんな人いたんだ。

942:デフォルトの名無しさん
09/01/15 16:29:33
英語的に　クー　よりも　キュー　だと思わね？

943:デフォルトの名無しさん
09/01/15 16:56:19
>>941
異端だな

944:デフォルトの名無しさん
09/01/15 16:57:46
どっちでもいいんじゃね

ちなみにnudeはニュードとヌードの両方の発音が有

945:デフォルトの名無しさん
09/01/15 17:13:01
キュもクも似てるし、まあ仮に間違って言ったとしても気づかれないんじゃない？

946:デフォルトの名無しさん
09/01/15 19:38:28
ホムペにはクーダと嫁
と書かれてたはず

947:デフォルトの名無しさん
09/01/15 22:24:41
"The cuda and my wife"

948:デフォルトの名無しさん
09/01/15 23:40:36
たしか英語版のfaqには
kuh-da
と書いてあったはず

949:デフォルトの名無しさん
09/01/15 23:53:57
2.1はJITサポートがミソ？

950:デフォルトの名無しさん
09/01/16 14:10:31
2.1をつついた人に質問
2.1ではカーネルに渡す引数でポインタ配列は使えますか?

951:デフォルトの名無しさん
09/01/16 18:48:04
共有メモリからグローバルメモリに連続領域をコピーするにはどうすればいいでしょうか
カーネルから呼び出せるmemcpyみたいなのがあれば教えてください

952:デフォルトの名無しさん
09/01/16 18:58:33
cudaMemcpyが長いようで、the launch timed out and was terminated.が
でてしまいす。これなんとかするほうほうありましたっけ?

953:デフォルトの名無しさん
09/01/16 21:53:12
>>952
転送量を減らす。

>>951
cudaMemcpy

954:デフォルトの名無しさん
09/01/17 00:16:47
>>951
スレッド内で連続的にメモリにアクセスするようなコードは望ましくない。
coalscedにアクセスするとなると、自分でコードを書くしかない。
共有メモリから転送を始める前に、同期を取ることも忘れずに。

>>952
参考までに、どんなハードウェアの組み合わせでどれくらい転送しようとしている?
8800GT辺りだと転送できるだけ全部転送してもタイムアウトしないのだけれど。

955:デフォルトの名無しさん
09/01/17 02:52:44
>>954
カーネル実行後の結果データコピーでエラーになっているのですが
sizeof(float)*1000バイトを100回連続でコピーしています。
100個とも別のアドレスに結果をコピーしたいためそうしてます。やはり
いっぺんにメインメモリにいったんコピーしてからやったほうがいいので
しょうか?

956:デフォルトの名無しさん
09/01/17 03:03:49
>>955
sizeof(float)*10000バイトを100回連続コピーでした。

957:デフォルトの名無しさん
09/01/17 07:16:40
>>954
＞ >>952
＞どんなハードウェアの組み合わせ

958:デフォルトの名無しさん
09/01/17 07:19:52
>>956
同一データ部分を100箇所にコピー？
独立の100箇所をコピー？

959:デフォルトの名無しさん
09/01/17 07:55:29
>>956
それ、本当にメモリのコピーでタイムアウトしてる？
カーネル実行後に同期取らずにすぐにメモリのコピー始めて
100回コピーのどこかでカーネル実行がタイムアウトしてるとか。

cudaMemcpyの戻り値は
Note that this function may also return error codes from previous, asynchronous launches.
ということだし。

960:954
09/01/17 08:44:25
>>955=956
えーと、都合数MB程度ってことか。それだったら>959の言うようにカーネル関数のタイムアウトじゃないかな。
DEBUGつきでCUT_SAFE_CALL(間違ってたら失敬)を使うか、自前でcudaThreadsSynchronize()を呼んで戻り値を見てみて。
カーネル関数起動後は、復帰を待たずにホストの処理に戻るからカーネル関数自体のエラーはカーネル関数起動では捕捉できないのよ。
で、まさかとは思うけど同一データを100回コピーならバス越しにやってはダメ。別のデータならホスト側の用件次第かな。
# でも多分、一旦ホスト側に転送してから分配した方が速い気がする。CPUとメモリ周りがよっぽど遅くなければ。

>>957
下手な突っ込みはお郷が知れるよ。この場合、CPUが遅いんじゃないかとかバスがx1なんじゃないかとかGPUがモニタ表示か何かで
よっぽど負荷が掛かっているんじゃないかとかGPU側メモリが巨大なSTELAなのかとか、組み合わせ次第で色々条件が変わってくるのよ。

961:デフォルトの名無しさん
09/01/17 09:52:40
>>960
＞下手な突っ込みはお郷が知れるよ。

質問者が答えてないのを指摘しただけなんだが…
誤解させてすまんかった

962:デフォルトの名無しさん
09/01/17 11:09:08
そういう意味か。あの書き方じゃぁ、誤解されるよ。
まぁ、大量転送じゃないと言う意味で、今回は関係なさそうだね。

963:デフォルトの名無しさん
09/01/17 15:38:06
>>952の話
>>957
ATIの統合チップセットでオンボードのATIビデオカードを無効
PCI-E 1.0 x16バスにGeforce9600GT 512MB
CPUはAMDデュアルプロセッサ(結構遅いやつ)

964:デフォルトの名無しさん
09/01/17 15:55:29
>>952の話
>>959
カーネル実行
a:cudaのエラーチェック(カーネルのエラーのありなし)
cudaThreadsSynchronize
for(100回
結果データをデバイスからホストにコピー
b:cudaのエラーチェック(直前のcudaMemcpyのエラーのありなし)
)
bのとこでthe launch timed out and was terminatedが出てるんです。

先週は時間がきてあきらめて帰ったのでここまでです。
言われたようにcudaThreadsSynchronizeの戻り値見てみたほうがよさそう
ですね。

965:デフォルトの名無しさん
09/01/17 18:45:09
>>951
共有メモリはカーネル起動時に動的に確保されるメモリ領域だから
カーネルが終了したら消えるし、１つの実行スレッドから全部コピーなんて論理的におかしいし
共有メモリはあくまで高速動作と同期が取れる作業用のメモリとして考えたほうがいい
共有メモリを使って計算した結果は１個づつ普通のグローバルメモリに入れてやるのが正しいやり方

>>952
10000を一度に転送して実行しても
1を10000回繰り返して転送しても
実行時間は大差ないんですよ
CUDAで実行する部分は出来るだけコンパクトにまとめて
呼び出す本体の方で特殊な演算関数くらいの感覚でループさせて使うのが正解

966:デフォルトの名無しさん
09/01/18 23:59:40
JCublasについて
おしえてエロイ人

967:デフォルトの名無しさん
09/01/19 11:04:50
>>925の話
cudaThreadsSynchronizeの戻り値をチェックしたら
the launch timed out and was terminatedが出ていました。結果コピーで
エラーで落ちてたのだけどエラーデータは、その前に実行していた
cudaThreadsSynchronizeの問題だったようです。

(cudaThreadsSynchronizeが正常になるまで待つとしても配列を
100x1000回、100x10000回、回すと1分待っても同期とれないようです。)

つまりカーネルに同期できない処理が書かれていたのが原因だと思います。
NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードで
かつそこのカーネルのdevPtrを参照するだけでなく書き戻す処理をや
っているのでそれがだめなのだと思います。恐らくこういうピッチ処理
の場合は、参照のみが許されるのでは?と思います。
問題のコードをこのpdfの変数に直して下記に書いておきます。
__global__ void myKernel(float* devPtr, int pitch)
{
for (int r = 0; r < height - 1; r++) {
float* row1 = (float*)((char*)devPtr + r * pitch);
float* row2 = (float*)((char*)devPtr + (r + 1) * pitch);
for (int c = 1; c < width - 1; c++) {
row2[c] = row1[c]とrow1[c + 1]とrow1[c -1]やらの計算
}
}
}

このp34-35を見て思うのはindex使わずにこんな処理書いてGPUでやる意味
あるの?と思いますが。自分で書いてなんなのですが

968:デフォルトの名無しさん
09/01/19 17:28:20
お前はあほかｗ

969:デフォルトの名無しさん
09/01/19 17:49:31
タイムアウトになる原因はそのでかいループのせい
せいぜいミリ秒単位でタイムアウトを判断してるから
ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目

cudaThreadsSynchronizeは発行したカーネルがすべて終了するのを待つだけ

グローバルメモリは読み書きは出来るが前後は保障されないので
１スレッドが書き込みする箇所は限定する必要がある

共有メモリを使って他のスレッドが書き込んだ内容を参照して利用する場合に
カーネル内部で___syncthreadを使う
これが本来の同期の意味

970:デフォルトの名無しさん
09/01/19 19:44:51
>>952の話
>>968のように言われるのは分かって書いてみたんだけど
NVIDIA_CUDA_Programming_Guide_2.0.pdfのp34-35のようなコードが
マニュアルに書いてあるのがおかしいと思う。
__global__ void myKernel(float* devPtr, int pitch){}
そもそもこんな書き方じたいが書けるけど間違えな使い方だと。
この書き方しているとこにやらないようにこの部分に×印つけてほしい。

あとはコンパイラがえらくなったらfor多重ループをうまく処理する
アセンブラぐらい作ってほしい。(OpenMPでパラレルfor指定すると#のタグだけ
でプロセッサ使ってきって高速化なるんだけどね。そこまでぐらい将来
的にはしてほしい。)

971:デフォルトの名無しさん
09/01/19 20:53:12
文句言っても何にもならない。
そう悟る時がくるまで気長に待ちましょうよ。
そうすれば成長しまっせ。

972:デフォルトの名無しさん
09/01/19 23:38:15
それよりも先ず、日本語を何とかしてくれ。
＞間違えな使い方
＞しているとこにやらないように
＞プロセッサ使ってきって

973:デフォルトの名無しさん
09/01/19 23:41:25
そのサンプルはあくまでもcudaMallocPitch()がどういう風にメモリを確保するかの説明だからなぁ。
きっとライターは、まさかそのまま動かそうとされるは思っちゃいまいよ。

974:デフォルトの名無しさん
09/01/20 02:02:49
ところでドライバ180なんだが2Dモードだかなんだかで消費電力抑える為に
ドライバが自動的にクロックを下げて3Dゲーム動作させると上がる仕様なんだけど
CUDA実行しても上がらないぞｗ

975:デフォルトの名無しさん
09/01/20 02:08:30
ドライバを替えてくださいw
つーか、それって169.09のときに指摘されて直したんじゃなかったのか?＞NVIDIA

976:デフォルトの名無しさん
09/01/20 03:40:32
d_a[0][blockIdx.x][blockIdx.y][threadIdx.x]
って使おうとしたら"expression must have arithmetic or enum type"
ってでたんだけど何がいけないんすか？
教えてください。

977:デフォルトの名無しさん
09/01/20 07:58:55
>>976
d_aがポインタなら、途中で解決できない状況があるのかも。

978:デフォルトの名無しさん
09/01/20 08:41:39
共有メモリで構造体を使いたいのだけど
例えば
struct data {
int a;
char b;
}

func<<<dim3(), dim3(), SIZE * sizeof(data)>>>();

__global__ kernel(){
__shared__ data d[SIZE];
}

こんな感じでやるとCUDAとC側のパック構造が違うせいだと思うけどおかしくなる
どうやれば出来るの？

979:デフォルトの名無しさん
09/01/20 08:59:56
パックが違うのだけが理由なら、
struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。

980:デフォルトの名無しさん
09/01/20 13:53:14
ただでさえ少ない共有メモリをそんな無駄に使えない

981:デフォルトの名無しさん
09/01/20 20:44:37
>>977
世の中夜勤帰りで朝から寝てる人だっているんだよ？
引っ越しの時ちゃんと挨拶行った？
顔合わせたら軽く会話するとかしてちゃんとコンタクト取り続けてる？
日頃からそういうコミニュケーションが取れてればいつ洗濯機を回していいのか
いつ静かにしなければならないのか
迷惑を掛けないように生活出来るはずなんだが

982:デフォルトの名無しさん
09/01/20 23:36:58
>>980
マジで言っているのなら、設計が悪い。
どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。

983:デフォルトの名無しさん
09/01/21 00:22:25
共有メモリが制限されてるのに無駄な領域作って
ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ

984:デフォルトの名無しさん
09/01/21 00:25:51
GPUのメモリレイテンシって１２とかの世界だぞ
CPU用のDDR2で５だからな
intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ

985:デフォルトの名無しさん
09/01/21 00:35:27
CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
coaxschedな読み書きができるなら、共有メモリより遅くないぞw

986:デフォルトの名無しさん
09/01/21 02:24:30
>>975
(エンバグの)歴史は繰り返す。

987:デフォルトの名無しさん
09/01/21 12:18:47
>>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
え？
共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど

ところでローカル変数をすべて共有メモリに置いた場合のが早いんだけど
ローカル変数ってデバイス上のメモリに確保されるだけで共有メモリより遅いレベルのメモリを使ってるっぽいね

988:デフォルトの名無しさん
09/01/21 18:11:15
あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。
レジスタみたいに使うのなら確かに関係なかったね。

で、レジスタよりも速いかどうかについてはptx見てみたいところ。

989:,,・´∀｀・,,）っ-○◎●
09/01/21 18:18:05
見た感じリードレイテンシはこれくらい

レジスタ＞＞Const Memory＞Shared Memory＞＞＞＞DRAM