- 1 名前：デフォルトの名無しさん mailto:sage [2007/09/17(月) 14:54:28 ]
- 本家
developer.nvidia.com/object/cuda.html
- 976 名前：デフォルトの名無しさん [2009/01/20(火) 03:40:32 ]
- d_a[0][blockIdx.x][blockIdx.y][threadIdx.x]
って使おうとしたら"expression must have arithmetic or enum type"
ってでたんだけど何がいけないんすか？
教えてください。
- 977 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 07:58:55 ]
- >>976
d_aがポインタなら、途中で解決できない状況があるのかも。
- 978 名前：デフォルトの名無しさん mailto:sage [2009/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 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 08:59:56 ]
- パックが違うのだけが理由なら、
struct data {int a; char b; char dummy[3];}とでもしておけばいいじゃん。
- 980 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 13:53:14 ]
- ただでさえ少ない共有メモリをそんな無駄に使えない
- 981 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 20:44:37 ]
- >>977
世の中夜勤帰りで朝から寝てる人だっているんだよ？
引っ越しの時ちゃんと挨拶行った？
顔合わせたら軽く会話するとかしてちゃんとコンタクト取り続けてる？
日頃からそういうコミニュケーションが取れてればいつ洗濯機を回していいのか
いつ静かにしなければならないのか
迷惑を掛けないように生活出来るはずなんだが
- 982 名前：デフォルトの名無しさん mailto:sage [2009/01/20(火) 23:36:58 ]
- >>980
マジで言っているのなら、設計が悪い。
どうしてもパディングしたくないくらい逼迫しているなら、int一個を切り分けて3バイトと1バイトで使え。
- 983 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:22:25 ]
- 共有メモリが制限されてるのに無駄な領域作って
ほとんどをグローバルメモリに追いやる方がよっぽど設計が悪いでしょｗ
- 984 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:25:51 ]
- GPUのメモリレイテンシって１２とかの世界だぞ
CPU用のDDR2で５だからな
intを内部でcharとして扱うプロセスを考慮しても共有メモリ使った方が早いんだよ
- 985 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 00:35:27 ]
- CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
coaxschedな読み書きができるなら、共有メモリより遅くないぞw
- 986 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 02:24:30 ]
- >>975
(エンバグの)歴史は繰り返す。
- 987 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 12:18:47 ]
- >>CUDAで共有メモリを使うこと自体、無駄な同期待ちが発生するから避けたいところだと思うが。
え？
共有メモリってCPUで言う所のただのレジスタみたいなもんで同期は指定しないと発生しないと思うけど
ところでローカル変数をすべて共有メモリに置いた場合のが早いんだけど
ローカル変数ってデバイス上のメモリに確保されるだけで共有メモリより遅いレベルのメモリを使ってるっぽいね
- 988 名前：デフォルトの名無しさん mailto:sage [2009/01/21(水) 18:11:15 ]
- あー、表現が悪かった。共有メモリを共有目的に使うと同期を取る必要が出てくるから同期待ちが発生するということ。
レジスタみたいに使うのなら確かに関係なかったね。
で、レジスタよりも速いかどうかについてはptx見てみたいところ。
- 989 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2009/01/21(水) 18:18:05 ]
- 見た感じリードレイテンシはこれくらい
レジスタ＞＞Const Memory＞Shared Memory＞＞＞＞DRAM
|
 |
