【GPGPU】くだすれCUDAスレ【NVIDIA】

1 名前：デフォルトの名無しさん mailto:sage [2008/03/22(土) 11:13:52 ] このスレッドは、他のスレッドでは書き込めない超低レベル、 もしくは質問者自身何が何だが分からない質問を勇気を持って書き込むスレッドです。 CUDA使いが優しくコメントを返しますが、 お礼はCUDAの布教と初心者の救済をお願いします。 CUDA・HomePage www.nvidia.com/cuda 関連スレ 【GPGPU】NVIDIA CUDA質問スレッド pc11.2ch.net/test/read.cgi/tech/1190008468/ GPUで汎用コンピューティングを行うスレ pc11.2ch.net/test/read.cgi/tech/1167989627/ GPGPU#2 pc11.2ch.net/test/read.cgi/tech/1188374938/ 659 名前：デフォルトの名無しさん mailto:sage [2009/05/11(月) 15:35:15 ] PyCUDAなんてあるのか、おもしろそう 660 名前：,,・´∀｀・,,）っ-○◎● mailto:sage [2009/05/11(月) 20:57:31 ] RuCUDAが必要だな 661 名前：デフォルトの名無しさん mailto:sage [2009/05/11(月) 21:41:12 ] >>656 参加したいGPUGRIDがどんなのか分からんが、 もし、倍精度浮動小数点の演算が必要なものなら、 GTX200シリーズじゃないと無理。 ちなみに、モニタがつながっているかPhysX指定がされてないと、 CUDAでデバイス列挙されないぽい。 662 名前：デフォルトの名無しさん mailto:sage [2009/05/12(火) 01:29:13 ] Py損とかルビィとか手続き型スクリプト言語は向かんだろ。 ocamlとか、Earlangとかの関数型言語をGPGPU対応に した方が御利益は大きいんじゃね? 並列計算の場合、副作用とか、計算の依存関係が有ると 性能出ないんで…。 663 名前：デフォルトの名無しさん mailto:sage [2009/05/12(火) 22:31:21 ] Earlang（笑） 664 名前：デフォルトの名無しさん mailto:sage [2009/05/16(土) 00:29:01 ] 統計解析ソフト「R」用のパッケージ「gputools」： ttp://cran.r-project.org/web/packages/gputools これのWindows版バイナリを作ってくださるネ申はいらっしゃいませんでしょうか。　m(ﾟ-ﾟ；)ｶﾐｻﾏ… 665 名前：デフォルトの名無しさん mailto:sage [2009/05/16(土) 10:27:22 ] >>660 こんなのはあるみたいですが ruby-opencl.rubyforge.org/ 666 名前：デフォルトの名無しさん mailto:sage [2009/05/18(月) 17:38:13 ] GPUの計算部分で レジスタを多く使っちゃうようにコンパイラで最適化されちゃうんだけど 部分的に無効にする方法はありませんか？ 667 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/05/18(月) 20:36:57 ] volatile 668 名前：デフォルトの名無しさん mailto:sage [2009/05/18(月) 23:05:56 ] >>667 ？volatileは最適化から外すだけで、レジスタには適応されるっしょ つか>>666 のレジスタ使ったら嬉しくない事ってのが想像できない。。。 669 名前：デフォルトの名無しさん mailto:sage [2009/05/18(月) 23:42:16 ] >>648 ＳＰだって年間何億個でも量産できるけど。 >>649 誰も、ＧＰＵを何個も用意しろとは言ってない。 670 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/05/18(月) 23:56:12 ] >>668　volatile __shared__ 671 名前：デフォルトの名無しさん mailto:sage [2009/05/19(火) 09:21:51 ] >666の動機によっては__shared__では何の解決にもならないような。 確認していないけれど、恐らくレジスタを使い回さずに消費しまくる方が速いんだろうねぇ。 672 名前：デフォルトの名無しさん mailto:sage [2009/05/26(火) 14:36:24 ] SP2＋未公開パッチで7RC以上に軽くなってるよ Windows Vista SP3 Part1 pc12.2ch.net/test/read.cgi/win/1241571715/225 673 名前：デフォルトの名無しさん mailto:sage [2009/05/26(火) 14:38:56 ] ｺﾞﾊﾞｰｸ 674 名前：デフォルトの名無しさん mailto:sage [2009/05/29(金) 18:11:27 ] SSE　４コアフルに使ったら 最上位品でも大差ないｗ 675 名前：デフォルトの名無しさん mailto:sage [2009/05/30(土) 15:06:57 ] code.google.com/p/thrust/ >Thrust is a CUDA library of parallel algorithms with an interface resembling the C++ Standard Template Library (STL). 676 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/05/30(土) 18:16:26 ] きた！STLきた！これで勝つる！　　 ないない 677 名前：デフォルトの名無しさん mailto:sage [2009/05/30(土) 21:28:10 ] brookみたいだな。 678 名前：デフォルトの名無しさん [2009/06/02(火) 17:20:15 ] 仮想マシン上でCUDAのインストールに成功した方はいらっしゃいますか？ 当方、ホストOS：Vista、ゲストOS：Ubuntu8.04です。 仮想マシンであるUbuntu上で、NVIDIAドライバ: NVIDIA-Linux-x86-180.22-pkg1.runを起動してみました。 すると、「 You do not appear to have an NVIDIA GPU supported by the 180.22 NVIDIA Linux graphics driver installed in this system」とエラーがでました。 GPUは、GeForce 9800 GTです。どなたか、アドバイスお願いします。 679 名前：デフォルトの名無しさん mailto:sage [2009/06/02(火) 17:50:02 ] 仮想マシンは無理 680 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 02:45:20 ] >>675 合計なんかCUDAでやって早くなるのか？ 681 名前：,,・´∀｀・,,）っ-○○○ mailto:sage [2009/06/03(水) 02:49:18 ] 分割統治法は並列化の基本だな 682 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 10:53:10 ] 合計求めるのは苦労したなぁ。 結局、分割数(128とか256とか)置きに足していって、その結果はPCで足した記憶がある。 683 名前：デフォルトの名無しさん mailto:sag [2009/06/03(水) 11:18:35 ] とりあえず公式の３つをインストールしたのですがTMPGEncで確認できませんみたいなことを言われました。 インストールするだけではcudaの恩恵を与れないのでしょうか? 684 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 11:31:01 ] すいません、直ぐ解決しましたorz ｸﾀﾞがちゃんと動いているか確認したいのですが方法はありますでしょうか？ 685 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 11:57:52 ] >>684 SDKをインストールしたのなら、サンプルもインストールしてビルドしてみよう。サンプルが動けば、大丈夫。 # 特に、deviceQueryは便利。 686 名前：デフォルトの名無しさん [2009/06/03(水) 14:12:00 ] こんにちは。CUDA初心者です。質問があります。 Visual C++ 2008、CUDA tool kit ver2.1、CUDA SDK ver2.1 で、サンプルのsimpleCUBLASをビルドすると、 1>LINK : fatal error LNK1181: 入力ファイル 'cutil32D.lib' を開けません。 と出ます。 そこで、CUDA SDKのlibを調べたところ、 cutil64D.libがあり、32のほうはありませんでした。 この場合、どうすればビルドできるのでしょうか？ 687 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 15:00:01 ] リリースバージョンをリンクする。 688 名前：デフォルトの名無しさん [2009/06/03(水) 16:01:01 ] リリース構成でビルドしたところ、今度は、 1>LINK : fatal error LNK1181: 入力ファイル 'cutil32.lib' を開けません。 と出ました。 CUDA SDKには、64があり、32はありません。 32と64の違いって一体何なのでしょうか・・・ 689 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 16:17:51 ] パスが通ってないんだろ 690 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 17:07:17 ] OSが64bitだと、32bitのCUDAライブラリはインストールされなかったような。 691 名前：デフォルトの名無しさん [2009/06/03(水) 17:07:40 ] ご回答ありがとうございます。 パスが通っていないということですが、 「パスを通す」について、詳しく説明していただけませんか？ 知識不足で申し訳ありません；； 692 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 17:43:34 ] 環境を名に使ってるかによるが、 Visual Studioだとプロジェクトのプロパティからインクルードするファイルがあるディレクトリのパスと、 libがあるディレクトリのパスをついかする linuxだとコンパイラのオプションに追加する 詳しくはぐぐれ 693 名前：デフォルトの名無しさん [2009/06/03(水) 18:04:37 ] 何度も回答していただきありがとうございます。 リンカの追加ライブラリを調べたところ、 ちゃんと、SDKのcommon/libが指定されてました。 しかし、この中には、cutil32.libではなく、64があります。 ということは、690さんのおっしゃるとおり、 32bitのCUDAライブラリはインストールされなかったということなのでしょうか？ もしそうでしたら、サンプルプログラムは64bitに対応してないが、 自分でプログラムを作る分には、上記のようなエラーはでないということでしょうか？ 694 名前：デフォルトの名無しさん mailto:sage [2009/06/03(水) 23:53:30 ] サンプルのリンカ設定を編集して64bitのライブラリをリンクすればいいんじゃね？ 695 名前：デフォルトの名無しさん [2009/06/04(木) 13:12:53 ] アクティブソリューションプラットフォームとプラットフォームをWin32からx64へ変更したところ、 エラーがなくなりました。 そのかわり、 1>------ ビルドのスキップ: プロジェクト: simpleCUBLAS ------ 1> ========== ビルド: 0 正常終了または最新の状態、0 失敗、1 スキップ ========== とスキップしてしまいました。 何が原因なのでしょうか・・・ 696 名前：デフォルトの名無しさん [2009/06/04(木) 18:26:30 ] threadIdx.xがうまく値を返さなく困っています。 最小のプログラミングだと　ちゃんとした値を確認できるんですけど、 規模のあるプログラム書いた物では、threadIdx.xをみると最大でも1000以内の数値が40000を超えていたりします。 かなりラフな書き方していて、グローバル変数使いまくってるのが意見ないのでしょうか？ __device__　 kouzoutai[2000];//グローバル変数 とか宣言しまくって搭載メモリーを超えちゃってるかもしれませんが、その時は明確なエラーとか出ますか？ 697 名前：デフォルトの名無しさん mailto:sage [2009/06/04(木) 19:59:22 ] エラーは基本的に出ない 698 名前：デフォルトの名無しさん mailto:sage [2009/06/05(金) 09:08:24 ] >>696 threadIdxはプログラムで変更できないと思うが。 どうやってthreadIdx.xを「みる|確認できる」としたのか知らんが、その手段を確認すべきでは? ptx出力を眺めればthreadIdx.xをコピーした先で壊してたりするのが分かるんじゃない? グローバル変数にたった2000要素くらいなんてことないと思うが、巨大な構造体でも置いてるの? あーそれから、メモリをオーバフローしてもコンパイルエラーも実行時エラーも出ないどころか、 突然システムごとフリーズするかもしれないから覚悟しておいてね。 699 名前：デフォルトの名無しさん [2009/06/06(土) 17:41:22 ] >>501 　　　？ 　　∧∧ 　 (´･ω) 　_｜⊃／(＿＿＿ ／ ヽ_(＿＿＿_／ ￣￣￣￣￣￣￣ 700 名前：デフォルトの名無しさん mailto:sage [2009/06/07(日) 01:25:14 ] 寝た子を起こしたなw 701 名前：デフォルトの名無しさん [2009/06/08(月) 08:06:10 ] >>699 　　！ 　 ∧∧ 　(･ω･ ) 　_| ⊃／(＿＿ ／ ヽ-(＿＿_／ ￣￣￣￣￣￣ 702 名前：デフォルトの名無しさん mailto:sage [2009/06/08(月) 20:11:26 ] 　 　　あ" 　 ∧∧ 　(◎ω◎) 　_| う／(＿＿ ／ ヽ-(＿＿_／ ￣￣￣￣￣￣ 703 名前：デフォルトの名無しさん [2009/06/09(火) 23:30:40 ] >>701 　　　　　ｵﾊﾖｰ!! 　　　　∧∧ ∩ 　　　 (`･ω･)/ 　　　⊂　　ノ 　　　　(つノ 　　　　 (ﾉ 　＿＿_／(＿＿_ ／　　(＿＿＿／ ￣￣￣￣￣￣ 704 名前：デフォルトの名無しさん mailto:sage [2009/06/10(水) 20:07:54 ] 質問。 グリッド・ブロック・スレッドのカーネル側での計算への応用方法が よくわからないのですが・・・・特に動きです。 const unsigned int tid =blockIdx.x * blockDim.x + threadIdx.x; がスレッドＩＤを示しているというのはわかるのですが・・・ for(int x=0;x<height;x++) for(int y=0;y<width;y++){ 　　　　　　　　sum=sum+tex2D(tex, x, y)*((float)(cos((M_PI*((x*width+y)+0.5)*tid)/number))); } outputdata[tid]=sum; の中で、tidはどういう働きをしてるんでしょうか？ 特に回収する配列であるoutputdataの動きが知りたいです。 わかる方ご教授お願いします・・・ 705 名前：デフォルトの名無しさん [2009/06/12(金) 04:58:52 ] >>703 　 ∧,,∧ 　( `･ω･)　ｳｰﾑ…ここは？ 　/　∽ | 　しー-Ｊ 706 名前：デフォルトの名無しさん mailto:sage [2009/06/12(金) 09:52:33 ] >>704 何が判らんのか判らん。 tidが個別のデータスレッドの番号を指すように働くことが判っているなら、 その番号でoutputdataにアクセスしていることくらい判りそうなもんだが。 仮に、blockDimが4でgridDimが2だとしたら、blockIdxは0か1、threadIdxは0から3になるわけで。 つまりはtidは0から7になるわけだな。 つーか、VIPのWikiにも情報があるよ。 vipprog.net/wiki/%E3%83%97%E3%83%AD%E3%82%B0%E3%83%A9%E3%83%9F%E3%83%B3%E3%82%B0%E8%A8%80%E8%AA%9E/CUDA.html 707 名前：デフォルトの名無しさん [2009/06/13(土) 00:18:28 ] >>730 カウンター3連打って言うけどスローで見るとカウンター2発に追撃1発って感じ まあそれでも凄いって言うか、何この異次元映像ｗ もはや芸術的とか圧倒的とかそういうレベルを超越していて笑いしか出てこないわけだが？ｗｗｗ しかもメディナは立ち上がるしｗｗ不死身かｗｗｗｗ この試合はボクシング以外の何か別の競技だわｗｗｗｗｗｗｗｗ 708 名前：デフォルトの名無しさん mailto:sage [2009/06/13(土) 00:24:10 ] 誤爆すまそ 俺が代わりに言っといてやったぜ 709 名前：デフォルトの名無しさん mailto:sage [2009/06/13(土) 01:35:36 ] CUBLASに加えてCULAPACKとか用意して欲しい LU分解もできない環境なんて 710 名前：デフォルトの名無しさん [2009/06/13(土) 04:07:09 ] GPGPUはDirectX11で花開くんだろ。 エンコード高速化対応お待ちしております。 DirectX 11講座（中編） 非ゲーマーのアナタのパソコンも DirectX 11演算シェーダーでスーパーコンピューターに変身する！ game.watch.impress.co.jp/docs/series/3dcg/20090327_79998.html GPGPUの場合は、組んだGPGPUプログラムがATIかNVIDIAのどちらかでしか動かない、 踏み絵的な「ATI対NVIDIA」なのだ。 これではGPGPUの進化、発展、普及が望めない。 そこでGPGPUの標準化プラットフォームが強く望まれるようになった。 このGPGPU標準プラットフォームを、Windows環境下で提供しようと出てきたのがDirectX 11ということになる。 DirectX 11は、このGPGPUの標準プラットフォームとして「DirectX Copmute Shader」（演算シェーダー）を提供する。 演算シェーダーでなにができるのか game.watch.impress.co.jp/img/gmw/docs/079/998/html/3dd1103.jpg.html 711 名前：デフォルトの名無しさん mailto:sage [2009/06/13(土) 08:42:36 ] >>710 エンコードって意外に並列処理できる箇所少ないんだけど。。。CUDAでどのベンダも対応してこないのはそのため。 Compute Shader出てきても状況は変わらないと気が 712 名前：デフォルトの名無しさん mailto:sage [2009/06/14(日) 15:19:59 ] それマルチ 713 名前：デフォルトの名無しさん mailto:sage [2009/06/14(日) 18:56:21 ] 今日本屋見てきたけど CUDAの入門書とかって無いね。 世の中にはあるのかな。 714 名前：デフォルトの名無しさん [2009/06/16(火) 16:56:23 ] __device__ int particle_position[30][30][50][30]; __device__ int particle_position_num[30][30][50]; 　： とか __device__ __constant__ float K_dWpress; 　： な感じで、結構たくさんグローバル変数をデバイス側に作ってるんだが、 これって、CUDA的にやっていい手法？てか数次元配列ってカーネル関数内で普通のcっぽくアクセスできる？ とりあえずメモリは足りるみたいだから、デバイスホスト間でのデータ転送は基本的にしなくて、 結果のみホスト側に持ってきたいんだけど。 コンパイルエラーはないけど、カーネル起動させたあとエラーが出て、 一瞬画面が消える。 715 名前：714 [2009/06/16(火) 17:22:05 ] 追記 ちなみに、Emuモードだと問題なく終了できます。 716 名前：デフォルトの名無しさん mailto:sage [2009/06/16(火) 17:29:37 ] >>714 CUDAはメモリ管理なんてしないから、デバイス側の変数がどれだけあるのか自分で把握すること。 一瞬画面が消えるのは表示用とCUDA用を同じGPUでやっている所為だと思われ。 もっと酷いときはフリーズするから要注意で。 717 名前：デフォルトの名無しさん [2009/06/16(火) 23:29:13 ] CUDA初心者です。 サンプルをいじくっているのですが、いくつかわからないことがあり質問させてください。 �@テクスチャメモリ キャッシュが効くからグローバルメモリよりも高速との事ですが、 時間を計ってみるとグローバルメモリと変わらない・・・。 どのような用途で使用するのでしょうか？ �A異なるカーネル関数の同時実行 ブロック、スレッド数を決めてカーネル関数を実行しますが、 cudaThreadSyncrinize関数やメモリコピー関数を呼ばなければ 同期待ちはしないですよね？ であれば、カーネル関数を2つ書いたときに同時実行するのでしょうか？ マニュアルを見る限りではできないようですが、会社の先輩が出来るというのでホントかなと。 お分かりの方宜しくお願いいたします。 718 名前：デフォルトの名無しさん mailto:sage [2009/06/16(火) 23:32:51 ] 俺も初心者だが。 1. 線形補完できるよ！ 速度面で違わないなら、それ以上はどう違うのか知らん。 2. なんか俺の動かしてる感触だと同時実行してる気がする。 少なくともCPUとGPUで同期はしていないよ。 719 名前：717 [2009/06/16(火) 23:59:21 ] >>718 サンクス。 �@は線形補間機能があるのは存じていますが、 補間せず生データを使用したいものでして・・・。 �Aは自分でも本家CUDAZONEで調べてみましたが、どうもできないっぽいです。 forums.nvidia.com/index.php?showtopic=84740&hl=cudaStreamSynchronize このスレの前にもありましたね。 よく見てませんでした、すみません。 cudaStream〜って何のためにあるんだか。 720 名前：デフォルトの名無しさん mailto:sage [2009/06/19(金) 00:40:24 ] CUDAを始めようかと考えていますが、下記の様な処理はCUDAで速くなりますか？ 721 名前：デフォルトの名無しさん mailto:sage [2009/06/19(金) 00:55:31 ] function hoge(const b:string):boolean; begin result:=true; end; 722 名前：デフォルトの名無しさん mailto:sage [2009/06/19(金) 00:58:07 ] >>720 すみません、間違えて書き込んでしまいました。 改めて、 typedef struct _table_t { float x,float y, float z; } table_t; typedef sturct _vec { float vector[4]; } vec; typedef union _vtemp_t { vec V[3], table_t data[4]; } vtemp_t; table_t DATA[100];(DATAにいろいろと情報をが入っている) vtemp_t temp; temp.data[0] = DATA[15]; temp.data[1] = DATA[43]; temp.data[2] = DATA[11]; temp.data[3] = DATA[80]; (tempのメンバdataにＤＡＴＡの値をランダムに代入) メンバdataがもつ、x､y､z情報をメンバVにx､y､z毎にまとめたいです。例えばこんな感じに。 V[0] = {data[0].x, data[1].x, data[2].x, data[3].x}; V[1] = {data[0].y, data[1].y, data[2].y, data[3].y}; V[2] = {data[0].z, data[1].z, data[2].z, data[3].z}; CellとかSSEならSIMDを使って出来ると思いますが、CUDAでも速くなるなら、 CUDAをやってみようかと思っています。 723 名前：デフォルトの名無しさん mailto:sage [2009/06/19(金) 01:25:58 ] それ計算じゃなくてただのメモリ転送じゃん。 CUDAでやることじゃないよ。 724 名前：デフォルトの名無しさん [2009/06/20(土) 14:41:38 ] Mac でCUDAをやろうとしています。 /Developer/CUDA/lib/にあるライブラリをライブラリのサーチパスに含めるには どの環境変数を設定すればよいのでしょうか？ LD_LIBRARY_PATHに追加してみたのですが、サーチしてくれません 725 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 14:43:15 ] >>724 DYLD_LIBRARY_PATH 726 名前：デフォルトの名無しさん [2009/06/20(土) 15:02:09 ] >>725 早速の回答ありがとうございます。 DYLD_LIBRARY_PATH に追加してみたのですがやはりうまくいきません。 ちなみに-Lで指定すればうまくいきます。 727 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 18:11:27 ] Coalesced、Non-Coalescedって簡単にいうとどういうこと？ 728 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 18:14:29 ] >>108 729 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 18:18:03 ] >>728 よくわからん 730 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 20:19:04 ] ttp://noridon.seesaa.net/article/105559613.html 731 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 21:42:57 ] 大前提として、ブロック内でデバイスメモリにアクセスするアドレスが連続していたら、コアレス。 んで対応アーキテクチャによってスレッド単位でアドレスが連続している必要があったりする。 まぁ簡単に言うと、バス幅有効活用できまっせ適な使い方をイメージすればいいかと。 なんでそのバス幅でアクセスできるアラインを考える事になる。 732 名前：デフォルトの名無しさん mailto:sage [2009/06/20(土) 23:08:40 ] >>726 実行時じゃなくて、ビルド時のリンクのこと？ DYLD_LIBRARY_PATHは関係ないから-Lでやってね。 733 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 01:03:25 ] >>726 LIBRARY_PATHでできなかったっけ？ 734 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 06:52:13 ] 配列で最初と最後の変数だけ別の処理をしたいのですが、 どのように書けばよいのでしょうか？ 従来のプログラムの場合、配列がa[N]だったばあい、 for(i=1;i<N-1;i++)と書けばよいのですが、 CUDAのカーネルで処理する場合どのようにすればよいのでしょうか？ 735 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 07:45:47 ] >>734 とりあえず、現状のカーネル書いてくれれば教えやすいのに 736 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 14:51:04 ] NVIDIAが配布しているということで、PhysXの質問をさせて下さい。 キネマティックなアクターが物体と接触した時に受ける力の取得方法が分かりません。 ご存知の方がいらっしゃいましたらご教示下さい。 737 名前：734 [2009/06/21(日) 17:50:38 ] >>735 例えば、今は __global__ void test(int thread_size, int one_size, float *res_d) {     int bx = blockIdx.x;     int tx = threadIdx.x;     int pos = bx * thread_size + tx;     if (pos==0)         res_d[pos] = cos(res_d[pos]);     else         res_d[pos] = sin(res_d[pos]); } このように、if文で0とN-1の場合だけ別の処理をするようにしているのですが、if文は遅くなるので、避けたいのですが。 具体的にやりたいのは、3次元の偏微分方程式を解く際の境界条件部分の計算です。 この場合各面、各辺、各頂点合わせてif文が20個以上になってしまいます。 何か良い方法はないでしょうか？ 738 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 18:12:51 ] >>737 最初と最後だけ別の処理なら、そこはCPUに任せて、 残りをGPUにやらせればいいと思うけど、それじゃだめなん？ posを+1して、並列数をN-2にすればいけそうだと思うけど。 あと、そのソースだと N-1 の処理が別になってない気がするのと、 posがN以上かどうかのチェックが無いのが気になる。 （ソースは単に省略しただけで、チェックはそもそも必要ないように組んでるのかも知れないが） 739 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 18:53:40 ] >>737 一般的にアセンブリレベルで分岐をなくしたいときは、こんなテクニックを使ったりする。 // i == Nのときだけsinをcosにする float s = sin(なんとか); floac c = cos(なんとか); int b = i == N; // true が (int)1に変換されることを期待する return (c * i) + (s * (1 - i)); もちろん3行目は実装依存なので、プロセッサのマニュアル見ながらいちいち確認はした方がいいんだが、 だいたいのプロセッサで使えるテクニック。 740 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 19:09:54 ] >>737 それだと全てのパターンでsinとcosの両方が計算されるから float s = sin(なんとか + b * PI / 2); にすれば？ 741 名前：740 mailto:sage [2009/06/21(日) 19:11:22 ] ごめん安価ミスった >>740は>>739へのレスです 742 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 19:37:25 ] CUDAにはsincosねーの？ また三項演算子でいいんじゃねーの？ 743 名前：デフォルトの名無しさん mailto:sage [2009/06/21(日) 23:01:37 ] >>737 あらかじめ、特殊条件について位相を補正しときゃいいんじゃないの？と思うのは俺だけ？ 744 名前：デフォルトの名無しさん mailto:sage [2009/06/22(月) 00:12:40 ] >>736 PhysXスレがゲーム製作技術にあるからそっち行け 745 名前：736 mailto:sage [2009/06/22(月) 01:52:13 ] >>744 誘導ありがとうございます 746 名前：デフォルトの名無しさん mailto:sage [2009/06/22(月) 10:53:08 ] >>742 あるよ。但し、ストリームプロセッサに一個しか超越関数演算機がないから使い方間違うと遅くなるけど。 747 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 08:26:25 ] >>746 超越関数演算機なんてハード的に存在したっけ？倍精度演算機じゃなくて？ 748 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 14:47:11 ] >>747 SFU(Super Function Unit)がある。 ただし、4SPで共有で、>>747のいうように8SPで共有するのは倍精度演算器。 sin()、cos()、tan()は確かに遅いけど、 sinとcosについては__sinf()と__cosf()を使えば、 精度は悪くなるけど1op/clockでできる。 除算以外の算術演算が8op/clockでできることを考えたらそれでも遅いことには変わりないけど。 749 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 15:22:26 ] 超越関数がレイテンシ1なんて凄いね SFUだけ100倍のクロック数とか してるのかなぁ 750 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 18:30:15 ] www.khronos.org/opencl/ OpenCL 1.0 リリース 751 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 18:36:49 ] First OpenCL demo on a GPU www.youtube.com/watch?v=r1sN1ELJfNo 752 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 18:43:26 ] www.nvidia.com/object/cuda_opencl.html ドライバ一式は登録しないといけないようだ 753 名前：デフォルトの名無しさん mailto:sage [2009/06/23(火) 21:47:02 ] >>748 SFUなんてあったのか、しらなかった。DPみたいに実行はSPと並行してできるのかな？ 754 名前：デフォルトの名無しさん mailto:sage [2009/06/24(水) 07:13:15 ] やっとFortran対応の話出てきたけど、有料なんかね PGI and NVIDIA Team To Deliver CUDA Fortran Compiler ttp://eu.st.com/stonline/stappl/cms/press/news/year2009/t2399.htm 755 名前：デフォルトの名無しさん mailto:sage [2009/06/24(水) 07:18:14 ] ていうかいつのまにかに2.3のベータ出てたのか 756 名前：デフォルトの名無しさん mailto:sage [2009/06/26(金) 12:01:06 ] GTX 285でようやくおれのようなへっぽこでも普通にcpuを凌駕するコードが書けるように なったようだね。 後は（値段はともかく）消費電力が下がってくれれば、大ブレイクしそう。 757 名前：デフォルトの名無しさん mailto:sage [2009/06/27(土) 12:19:52 ] ３００Wなんてクアッドコアマシンを４台くらいクラスタリングしたような消費電力だからな 758 名前：デフォルトの名無しさん mailto:sage [2009/07/01(水) 03:13:56 ] >おれのようなへっぽこでも普通にcpuを凌駕する だったらCPU要らんがな。 むしろCPU<->GPU間の通信がボトルネックなので、CPU無くして GPUに直接I/O繋げられるようにしろや！ 759 名前：デフォルトの名無しさん mailto:sage [2009/07/01(水) 21:14:50 ] >>758 それなんていうTesla?

---

---

[ 続きを読む ] / [ 携帯版 ]

read.cgi ver5.27 [feat.BBS2 +1.6] / e.0.2 (02/09/03) / eucaly.net products.
担当:undef