【GPGPU】NVIDIA CUDA ..
790:デフォルトの名無しさん
08/11/27 00:32:04
そいえばCUDAって1つのカーネルのサイズが制限されてない?
でかいやつがまったく動かなくて苦労したんだけど
791:デフォルトの名無しさん
08/11/27 01:39:04
どの位かは知らんが、そりゃぁ制限はあるだろうねぇ。
792:,,・´∀`・,,)っ-○◎●
08/11/27 01:54:09
64Kのコンスタントメモリがあるじゃん。
コンスタントメモリは自分自身では中身の入れ替えは不可能。
あとはわかるよな?
793:デフォルトの名無しさん
08/11/27 02:23:46
cudaで自己書き換えプログラムってできますか?
794:,,・´∀`・,,)っ-○◎●
08/11/27 02:37:41
GPUのカーネルコード自身で書き換えるって意味なら無理。
PTXのバイナリコードを動的生成とかなら何かやれば可能かもしれない。
その辺の資料を中の人に要求したら
「機密事項ですのでお答えできません」
795:デフォルトの名無しさん
08/11/27 17:49:14
大学の研究室にCUDAプログラミング用のコンピュータが導入された!と喜んでいたら、
HP ML115 + げふぉ8400GSカードだった…orz 学習用仕様で萎えた…
796:デフォルトの名無しさん
08/11/27 17:56:05
まぁ、学部生なら十分だろ。
どうせ大した論文も書かないくせに。
797:デフォルトの名無しさん
08/11/27 18:59:33
>>796
お前みたいなの、必ず湧くよなw
人を馬鹿にしたら、自分が偉くなるとでも思ってんの?
798:デフォルトの名無しさん
08/11/27 19:42:03
でも俺もそこまでは言わないにしても
学習用仕様で萎えたってセリフは贅沢だと思うよ。
799:名無し募集中。。。
08/11/27 19:46:05
値段を見て萎えたって意味だろうが性能的には十分な気がするんだけど
800:デフォルトの名無しさん
08/11/27 19:54:22
そもそもGPGPGUはなんなんだ?8800はVGA用途だろ
801:デフォルトの名無しさん
08/11/27 20:23:28
>>797
この程度で反応するなよw
社会でやっていけないぞ
802:デフォルトの名無しさん
08/11/27 22:49:14
>>795
それなんて俺
CUDA使ってみたくて買い換えを考えていたけど
PCI Express x16バス搭載で1万だったのでML115と
玄人志向の8400GSカード買ってきてCUDA環境を手に入れたぜ
メインで使ってるマシンより性能がよくてこのままメインになりそうな予感
803:デフォルトの名無しさん
08/11/27 22:51:16
事前に断っておくと俺は批判されるとチンコがたつよ
804:デフォルトの名無しさん
08/11/27 23:04:29
自宅でCUDA使ってる人のコンピュータスペックってどんなもん?
805:デフォルトの名無しさん
08/11/27 23:05:58
>>804
c2d e6320
806:デフォルトの名無しさん
08/11/27 23:11:07
>>804
PenD920 GeForce9600GT
807:デフォルトの名無しさん
08/11/27 23:31:34
>>804
Core i7 920 + GeForce GTX 280
808:デフォルトの名無しさん
08/11/28 00:15:30
>>804
Xeon + GeForce GTX280
809:798
08/11/28 03:29:12
>>799
いや、新技術半可通の俺が察するにGeF8000系は
CUDA対応ハードの底辺だから文句言ってるんだよ。
研究室で導入する(自腹じゃない)んだからもっと良いのよこせってことだろ?
ぶっちゃけ研究といっても今の時期じゃ全体的に未成熟だし
そのスペックでさえ限界動作なんぞせんだろと思って贅沢だなと。
810:デフォルトの名無しさん
08/11/28 03:39:20
>>800
そんなことはない。GPGPUは対応ハードならグラボ1枚でも2枚でも一応使えるんだよ。
実例は押さえてないんでなんともいえないが
競合さえおきなきゃRadeon(VGA)・GeF8800以上(GPGPU専用)ということも出来るはず。
用途と価格の兼ね合いもあるだろうけど、
とりあえず試したいとかコストパフォーマンス的に有利な部分もあるかもしれん。
4gamerの記事が参考になるよ。
URLリンク(www.4gamer.net)
811:デフォルトの名無しさん
08/11/28 04:14:57
>>809
底辺どころかメインストリーム
812:798
08/11/28 06:47:41
>>811
ラインナップ上は確かGeF8世代からの対応だから一応底辺だろう。
>>795の萎える理由に対しての推測だからこの説明でいいんじゃね?
もっとも現状ではSP数の差による性能差がそこまで決定的じゃないから
2x0とかQuadroを期待してるんだったら贅沢なんじゃねって話。
813:デフォルトの名無しさん
08/11/28 07:50:33
GT2x0:性能は兎も角、(電源などの)要求仕様が一般的でない。
QuadroFX:同程度の性能のGeforceの数倍の価格と言う段階で、論外。
# なんて書くと、「アキバ的発想云々」って言われるんだけどねw
もしこれらが理由なら、見識不足と言わざるを得ない。
単に、>795はML115が気に入らないんだろw
まぁ、今時なら8400GSでも8600GTでも9600GTでも大して値段変わらんと思うが。
814:デフォルトの名無しさん
08/11/28 11:07:07
ML115の電源じゃパフォーマンスレンジのGPUは無理ね
815:名無し募集中。。。
08/11/28 11:14:47
ML115の電源はサーバー向けだけあってそれなりに良い電源だよ
電源投入後の全開爆音は凄いけど
ちなみに容量は370W
816:デフォルトの名無しさん
08/11/28 11:52:57
それじゃ精々8600GTSか9600GT止まり。8400GSは案外無難な選択じゃないのかな。
817:デフォルトの名無しさん
08/11/28 13:12:49
ちょっと誰か日本の公式フォーラムのcudaggさんの日本語をなんとかしてあげて
818:デフォルトの名無しさん
08/11/28 13:51:35
Ubuntuの話なら、一応意味は判るでしょ。
つーか、ここでフォローしても意味ないしw
819:デフォルトの名無しさん
08/11/28 18:56:55
>>813
GT2x0で一般的でないとか言ってたらCUDA自体一般的じゃないだろjk
ML115が不服と言うのでも十分贅沢……。俺が行ってた某理系大学なんか(ry
820:流石理系大学出は日本語読むことさえ放棄しているようだな。
08/11/28 18:59:40
>>819
GT280は200Wの電源とそれに見合う通風環境が必要になるってこった。
この際、CUDAが一般的かどうかが問題じゃない。
821:考えること放棄してるやつよりマシだよ。
08/11/28 19:53:01
>>820
CUDAスレでの話だからマシン云々というよりは
単に出たばっかのGT2x0系じゃなかったことに萎えてるのかと思っただけじゃん。
そもそも研究室とかだったら要求仕様が一般的でないなんてことは些細なことだ。
例えばOracle高価だから絶対使いませんなんて真似しないだろ。
そういうところは必要あれば用意するだろうよ。
GT2x0載せるとしたらそりゃML115じゃ無理だろうけど、それに載せろとは誰も言ってないし。
一般的云々言われなきゃ俺だってCUDAが一般的かどうかなぞ持ち出さんわ。
822:デフォルトの名無しさん
08/11/28 19:55:38
グラボの方に萎えてると思ったもうひとつの理由が、
CUDAプログラミング用ということだから
マシン自体のスペックはそれほど高くなくても良いはずなんだ。
電源さえ足りてればな。深読みしすぎた俺が悪かったよ……。
823:デフォルトの名無しさん
08/11/29 01:30:32
そこいらでレスをつけてるやつらの中に、8000系では一部使えない命令がある件を指摘するやつが居ないのは何でだぜ?
お前ら実は妄想だらけでなんもしてないんじゃないの?
824:デフォルトの名無しさん
08/11/29 01:46:30
8400GSならCompute Capability 1.1だから
GT200で追加された命令でもなければ使えるわけだが。
825:デフォルトの名無しさん
08/11/29 02:59:51
>>823
1.0世代は8800GTX,GTS,ULTRAに積まれたG80くらいで、8600GTや8400GSに積まれたG84やG86は
1.1世代だと言う知識もなしにここ見てたの? 他人を妄想だらけなんて指摘してられる状況じゃないじゃんw
826:デフォルトの名無しさん
08/11/29 03:20:49
CUDAを単純に使いたいやつは、
HP ML115 (URLリンク(nttxstore.jp))
と
G98の8400GS買えばおkだよ。
827:デフォルトの名無しさん
08/11/29 19:48:41
グラボ早い順に並べるとどんな感じ?
やっぱりSP数×クロックなの?
828:デフォルトの名無しさん
08/11/29 21:13:40
仕様を限定せずに一般的に言うなら、そりゃそうだ。他にどんな要素が来ると?
特定の応用でメモリ転送が重視されるのならメモリ帯域やバス仕様にも依存するし、
メモリ量が多い方が高速なアルゴリズムを採用できる応用ならメモリ量にも依存するわけだけど。
あー、1.0世代だとストリームが使えないと言う大きな欠点はあるね。
829:デフォルトの名無しさん
08/11/30 00:11:00
チップ外は考えてなかったφ(..)フムフム…
後からSPが減ってクロックが上がった製品が出るじゃない?
だいたい値段は一緒だけど、劇的に性能が違うのかなと…
830:デフォルトの名無しさん
08/11/30 01:46:04
同じ名前のSP減ったのは、恐らく歩留まりの関係で在庫処分しているだけだと思う。
使い方にも拠るけど、クロックが利く応用なら結構変わるかもね。
# SP数が利く応用なら逆に遅くなりかねない。
831:デフォルトの名無しさん
08/12/01 03:44:29
やっぱみんなPDEとか立てたり解けたりできるの?
832:デフォルトの名無しさん
08/12/01 06:49:36
そうだね
SDEとかね
833:デフォルトの名無しさん
08/12/01 23:36:29
>>795
あぅっ、それはウチかも知れない。
まぁどーせ学生の作るプログラムなんて、レポート見てると
自分で作ってるのは15%ぐらいで、残りの人は友人のコピペ
改変か、ググって見付けたページのコピペ改変だからな。
予算の都合ってやつで、1人1台用意するならML115+GF8400
に成ってしまう。許せ。
おっ、これは?!と思うプログラムが出てきたら、GTX280で
走らせてやるから頑がってくれ。
834:デフォルトの名無しさん
08/12/01 23:40:13
>>816
PCI-Ex補助電源コネクタ無いので、GF8600GT
か、GF9500GTまでが無難なところ。
835:デフォルトの名無しさん
08/12/02 00:15:15
texture<float, 1, cudaReadModeElementType> tex;
のテンプレートで
> error C2018: 文字 '0x40' は認識できません。
ってエラーが大量に出るんだが、みんな出ないですか。
オレッスカ
textureのテンプレート使わなければでないのですが。
CUDA 2.1 + VC++ Express 2005
836:デフォルトの名無しさん
08/12/02 00:20:06
>>835
PDFからコピペしてない?
837:デフォルトの名無しさん
08/12/02 00:36:23
>>836
一回行コピペしてエラー出たので消して書いたのですが、
ファイルに変なコードが残るとかありますかね。
でもその行消すとエラーなしで、手書きで書いたら、
↓のようになります。
1>.cu(54) : error C2018: 文字 '0x40' は認識できません。
...
1>.cu(54) : error C2018: 文字 '0x40' は認識できません。
1>.cu(54) : error C2065: 'COMPILER' : 定義されていない識別子です。
1>.cu(54) : error C2146: 構文エラー : ';' が、識別子 'ERROR' の前に必要です。
1>.cu(54) : error C2065: 'ERROR' : 定義されていない識別子です。
1>.cu(54) : error C2143: 構文エラー : ';' が '<template-id>' の前にありません。
1>.cu(54) : error C2146: 構文エラー : ';' が、識別子 'tex' の前に必要です。
1>.cu(54) : error C2275: 'texture<T,dim,__formal>' : この型は演算子として使用できません
1> with
1> [
1> T=float,
1> dim=1,
1> __formal=cudaReadModeElementType
1> ]
1>nv_som_gpu.cu(54) : error C2065: 'tex' : 定義されていない識別子です。
838:837
08/12/02 00:46:22
texture<float1, 1, cudaReadModeElementType> *tex = new texture<float1, 1, cudaReadModeElementType>;
ならOKなので、スタックにtextureを置くとダメな条件があるのでしょうか。
839:837
08/12/02 01:31:47
kernelに引数で渡せないので
スタックにおいたらダメみたいでした。
解決。
840:デフォルトの名無しさん
08/12/02 23:13:52
まだ、初めてみようかと考えている初心者以前のものです。
現状の自作数値解析プログラムはMPIで各ノードの各コアにプロセスを振っております(ノード内もMPIで並列)。
これの自然な拡張としては、MPIの各プロセスがGPGPUを使うという形になるかと思います。
しかし、マルチCPUのノードの場合、ひとつのGPUを共有することになります。
一つのプロセスがGPUのシェーダのうち1/4だけ使って、4プロセスから同時にGPUを使うなんてことは可能なのでしょうか?
841:デフォルトの名無しさん
08/12/02 23:47:13
>>839
テクスチャはコンパイルするとただの数値みたいになるので、
グローバル変数としてそのまま使う以外のことはほとんど不可能、らしいよ。
関数にパラメータとしてを渡すのも、ポインタを得るのも無理。
842:デフォルトの名無しさん
08/12/03 00:02:52
>>840
CUDAを使う場合、4プロセスから同時に使うことも可能。
但し、GPUをどう割り振るかはCUDAドライバの御心次第。
厳密に制御するには、GPU1枚ごとに担当スレッドを設けることになる。
# その場合、プロセス間通信でJOB型にするか1プロセスだけでGPUを占有するかは設計次第。
それはいいけど、プロセス全部分けるとマルチスレッドに較べて効率落ちないかい?
843:840
08/12/03 00:20:11
>>842
ありがとうございます。
やはりリソースの競合が起きるようですね。
マルチスレッドは単純に覚えることが増えるのでやってませんw
スパコンのマニュアルでMPIを覚えたので、WSクラスタでもそのままの手法を持ってきてます。
一応、ノード内では共有メモリを使って通信するようなオプションでmpichをインスコしてるので多少はマシでしょうという感覚です。
844:名無し募集中。。。
08/12/03 02:19:05
NVIDIA、PhysX/CUDAを活用する「パワーパック」の第2弾を提供開始
URLリンク(journal.mycom.co.jp)
845:デフォルトの名無しさん
08/12/03 03:31:23
0x40でググったら
>文字 '0x40' は認識できません。 原因: ソースコード中に全角空白 ' ' が使われています。
まさか、な
846:デフォルトの名無しさん
08/12/04 01:11:01
こんな感じで網羅的に点を回転させるプログラムを書いているのですが、
Thread数が320を越えるあたりでrotate関数に全ての引数が渡らなくなってしまいます。
その時でも計算量が少ないからか、roll_axisのz成分などは、引数として機能しています。
roll_axisのx,y成分も適当な値に変えると引数として働くようになります。
rotate(&coord, &pitch_axis,ANGLE);をコメントアウトすれば,roll_axisのx,yはThread数が多くても引数として渡ります。
計算量が多くなると(特に三角関数?)起きる気がするのですが、何が回避する方法はあるのでしょうか?
--ptxas-options=-vはUsed 42 registers, 56+28 bytes lmem, 2080+32 bytes smem, 8024 bytes cmem[0], 88 bytes cmem[1]て出ます。
-deviceemuでは正常に動作します。
847:846
08/12/04 01:13:18
>>846のソースコードです。
####kernel(Thread1つが回転させる点1つに対応)####
__device__ runDevice〜〜の一部
float4 yaw_axis = make_float4(0, 0, -1, 0);
for(iyaw = 0; iyaw < limY; iyaw++){
float4 pitch_axis = make_float4(-sin(radian*iyaw), cos(radian*iyaw), 0, 0);
for(ipitch = 0; ipitch < limP; ipitch++){
float4 roll_axis = make_float4(cos(radian * iyaw) * cos(radian * ipitch), sin(radian * iyaw) * cos(radian * ipitch), sin(radian * ipitch), 0);
for(iroll = 0; iroll < limR; iroll++){
if(tid == iroll + ipitch * limR + iyaw * limR *limP)
g_mem[tid] = coord.x; //回転できているか確かめる。
rotate(&coord, &roll_axis, ANGLE); //回転させる点の座標と、回転軸と角度を与える。
}
rotate(&coord, &pitch_axis, ANGLE);
}
rotate(&coord, &yaw_axis, ANGLE);
}
__device__ void rotate(float4 *coord, float4 *axis, int angle){
coord->x = axis->x;
coord->y = axis->y; //とりあえず現段階では引数が渡るか確かめてるだけ
coord->z = axis->z;
}
848:デフォルトの名無しさん
08/12/05 14:25:22
誰か、すげぇ簡単単純だけど、真理を付いてる様な神サンプル晒してよ
849:デフォルトの名無しさん
08/12/05 15:02:45
>>848
どんなの? SDKのサンプルやbioのblog辺りじゃお気に召さない?
具体的なテーマがあって、実装が難しくなさそうなら作ってもいいけどね。
850:デフォルトの名無しさん
08/12/05 21:35:18
ATI Streamから来ました。
スイマセン、場違いな場所に来てしまった。
AもまだなのにCなんて出来ません
851:デフォルトの名無しさん
08/12/05 23:37:57
共産主義者の書いたマンガによると、Aより簡単らしいぞ。
852:デフォルトの名無しさん
08/12/06 00:23:10
最近の若い連中はしょっぱながCだからな
853:デフォルトの名無しさん
08/12/06 19:46:07
cudaってさ
C-daにしておけばギャグっぽくてよかったんじゃね?
854:デフォルトの名無しさん
08/12/07 12:03:03
それならCarracudaの方がw
855:デフォルトの名無しさん
08/12/08 22:50:38
teslaの話題が少なくて絶望した!
856:デフォルトの名無しさん
08/12/08 23:03:23
えー、QuadroFX5600のアナログ回路をとっ外しただけの代物の、何を語れと言うのさ。
聞いてくれたら答えるけど。
あ、4桁シリーズは白根。どうせ事情は一緒だと思うけど。
857:デフォルトの名無しさん
08/12/08 23:44:28
>>855-856
このスレの住人は GPGPUの用途にQuadroFX使ってるの・・・?
ていうかそもそもみんな何使っているの?
858:デフォルトの名無しさん
08/12/08 23:58:17
ML115最強伝説
859:デフォルトの名無しさん
08/12/09 00:12:18
Tesla D870 が至高
860:デフォルトの名無しさん
08/12/09 00:59:33
>>857
このスレともう一つのCUDAスレを見ると、QuadroFXのメリットが書かれていると思うが。
# 普通、使うわけないだろって。
8800GTX、QuadroFX5600、Tesla C870の価格を見る限り、Teslaは未だましだと思う。
>>859
D870ってNVIDIA謹製ミニタワーにC870を二枚入れた代物だっけ?
あれだったら同じ筐体のQuadroPlexの方が潰しが利くと思うのだけど。
861:デフォルトの名無しさん
08/12/09 01:51:03
GeForce GTX 280 極上
862:デフォルトの名無しさん
08/12/09 20:16:56
9600GT買ったらCUDAでエラー出まくりで問い合わせたら
CUDAのようなメーカー付属のドライバで対応してない機能は保障外とのこと
CUDAをやるために買ったのにCUDAが動くことを保障してないなんて
はっきり言って詐欺だよ
863:デフォルトの名無しさん
08/12/09 20:51:59
DELLにVisual Basicで組んだプログラムでエラー出るんだけど?
って聞いてるようなもんだな。
864:デフォルトの名無しさん
08/12/09 20:54:42
>>862
nvidiaのドライバだとどうなんだ?
マジに教えてくれ
865:デフォルトの名無しさん
08/12/09 21:07:30
>>863
CUDA SDKのテストプログラムでエラーが出まくる
実行する度に計算結果が違ってくる
>>862
どういう意味?NVIDIAのドライバ以外にCUDAが動くの?
付属のCDのドライバ以外はNVIDIAの公式ドライバでも保障しないと言われた
866:デフォルトの名無しさん
08/12/09 23:35:15
>>862
一体どこの糞メーカのPCなんだ?
867:デフォルトの名無しさん
08/12/10 23:15:44
>>865
>付属のCDのドライバ以外はNVIDIAの公式ドライバでも保障しないと言われた
それが普通ですよ。販売メーカとしては、自分のところでテストした
ドライバ以外は保証してないよ。
だから公式ドライバ入れる時とかもよく言われるでしょ。
トラブっても自己責任で、ってね。
868:デフォルトの名無しさん
08/12/11 00:00:11
まあ当たり前なのは分かるんだけど、そこんとこ盲点だよ。
付属のバージョンだと根本的にCUDA機能が入ってないからね。
必然的にアップデートしないといけないのに保障外でしょ。
CUDAは素晴らしいとか大々的に宣伝してるけど動くかどうかは運次第という対応だからね。
これが世間に知れたらCUDAなんて誰も相手にしなくなると思うけど。
869:デフォルトの名無しさん
08/12/11 00:16:59
tesla売るためだろ
870:デフォルトの名無しさん
08/12/11 00:19:56
そうだろうね。TESLAは唯一保障された製品だからCUDAやりたい奴はTESLA買えってことなんだろう。
一般人には何の恩恵もない誇大広告だったというわけですか。
871:デフォルトの名無しさん
08/12/11 00:23:48
メーカー純正フロッピーディスク以外使わない奴なんていなかっただろ。
みんなそうやって保証外の物を使ってきているんだ。今更何を言っているんだか。
872:デフォルトの名無しさん
08/12/11 00:27:34
で、そんなふざけた事言うのはどこなんだ?
873:デフォルトの名無しさん
08/12/11 00:27:39
どこかの誰かが同様にCUDAが使えないとなると自分がCUDAを使う意味がない。
だからCUDAは必要なくなるので別に何も問題ない。
誰もが使えないプログラムなんて作っても時間の無駄でしょ。
保障がないってのはそういう事よ。個人的な恨みとかじゃなく。
874:デフォルトの名無しさん
08/12/11 01:57:49
ノートPCなんてNVIDIAが公式ドライバを提供しないんだぜ。
メーカーも提供しないから、アマチュアの作ったものを探してくるしかないっていう。
875:デフォルトの名無しさん
08/12/11 04:03:59
あー、WindowsでCUDA試してないな、うちのMac
876:デフォルトの名無しさん
08/12/11 21:46:41
保証が無いからどうだとかいうのは自作板でやってくれ。
877:デフォルトの名無しさん
08/12/12 02:15:59
>>857
犬板にも書いたけど、JetWay HA05-GTのオンボッボで動きますた。
USBメモリ起動で、薄いCUDA nodeの出来上がりです。
オンボッボももっと沢山SPU載る日はいつかなぁ。
878:デフォルトの名無しさん
08/12/12 18:46:10
【GPGPU】CUDA/ATI STREAM 速度・画質検証スレ
スレリンク(jisaku板)
879:デフォルトの名無しさん
08/12/12 20:35:23
これからCUDAを使おうと思っているのですが、visualC++でうまくビルドできません。どなたか解決方法を教えてください。
環境はXPx64Edition,QuadroFX570でx64版のバージョン2.0のドライバとToolkitとSDKを入れました。
VisualC++は2005Expressです。
880:デフォルトの名無しさん
08/12/12 20:39:35
すみません、879ですが、ビルドできないのはSDKの中にあるprojectsのsimpleTemplateです。
webで調べてみると皆さん問題なく実行できるようなのですが。。。
881:デフォルトの名無しさん
08/12/12 22:05:56
どんなエラーが出るとかの情報はないの?
882:879
08/12/12 23:26:03
失礼いたしました。下のようなエラーがでます。
nvcc fatal : Visual Studio configuration file '(null)' could not be found for installation at 'C:/Program Files (x86)/Microsoft Visual Studio 8/VC/bin/../..'
>>444さんと同じような問題と思われます。VisualC++のコンパイラとの相性が悪いのでしょうか。詳しい方、よろしくお願いいたします。
883:デフォルトの名無しさん
08/12/13 06:12:02
こんな連中ばっかりだし、クマには生きづらい世の中になったぜ
884:デフォルトの名無しさん
08/12/13 12:41:18
C++もろくに使えない人がいきなりCUDAに手を出すのもどうか
885:デフォルトの名無しさん
08/12/13 14:55:07
なんか64bit環境でもVCのclは32bit版使わないとだめみたいなことCUDAフォーラム(英語のほう)に書いてあったよーな
886:デフォルトの名無しさん
08/12/13 14:55:41
nvccから呼ばれるほうのclのことね>32bit版
887:デフォルトの名無しさん
08/12/14 01:21:53
>>882
パスが通ってないんじゃね?
888:879
08/12/17 18:34:27
PATHは通っていると思います。消してみると、
nvcc fatal : Cannot find compiler 'cl.exe' in PATH
となりました。platformSDKのclにすると↓のようになりました。
nvcc fatal : nvcc cannot find a supported cl version. Only MSVC 7.1 and MSVC 8.0 are supported
結局、CUDAtoolkitを32bitにしましたところOKとなりました。885さん、886さん、887さんありがとうございました。
889:デフォルトの名無しさん
08/12/17 20:02:14
最適化してregisterの数を減らしたいんですが、
なんか心がけることとかコツってありますか?
890:デフォルトの名無しさん
08/12/17 23:45:04
ptx出力を読め。以上。
891:デフォルトの名無しさん
08/12/18 00:41:24
なんでわからないのでご教授を・・・・
VS 2005にカスタムウィザードを組み込んでパスも通して空の.cuファイルに書き込んだのに
1>LINK : fatal error LNK1181: 入力ファイル '.\Debug\sample.obj' を開けません。
と出るのですが・・・・何か環境設定間違っているのでしょうか?
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
シェーダプロセッサってエラーがあってもゲームでは無視されて
認識出来ないレベルで画面が崩れるだけだから
一見正常に動いてると思っててもCUDAは動かない人が多いかも
ドライバが対応してるとか以前にハードが用件を満たしてない
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を経由しなきゃ目いっぱい使えるかもね。
ただ、スレッドインターリーブすると1スレッドあたりで使えるレジスタ本数が減っちゃうんだよね。
メモリレイテンシを隠蔽するならスレッドを目いっぱい使ったほうがいいし
逆に一時変数を何度も再利用する場合は、thread/blockを減らして1スレッドあたりの仕事量を増やしたほうがいい。
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
共有メモリはカーネル起動時に動的に確保されるメモリ領域だから
カーネルが終了したら消えるし、1つの実行スレッドから全部コピーなんて論理的におかしいし
共有メモリはあくまで高速動作と同期が取れる作業用のメモリとして考えたほうがいい
共有メモリを使って計算した結果は1個づつ普通のグローバルメモリに入れてやるのが正しいやり方
>>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
お前はあほかw
969:デフォルトの名無しさん
09/01/19 17:49:31
タイムアウトになる原因はそのでかいループのせい
せいぜいミリ秒単位でタイムアウトを判断してるから
ミリ秒以下で応答しないようなカーネルは絶対に書いては駄目
cudaThreadsSynchronizeは発行したカーネルがすべて終了するのを待つだけ
グローバルメモリは読み書きは出来るが前後は保障されないので
1スレッドが書き込みする箇所は限定する必要がある
共有メモリを使って他のスレッドが書き込んだ内容を参照して利用する場合に
カーネル内部で___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実行しても上がらないぞw
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];}とでもしておけばいいじゃん。
次ページ最新レス表示スレッドの検索類似スレ一覧話題のニュースおまかせリスト▼オプションを表示暇つぶし2ch
4327日前に更新/252 KB
担当:undef