CUDAの画像処理はテクスチャが必須ですな
関連記事
- エヌビディアの時価総額300兆円超え AIブームうけ最速で達成 【2024年02月24日(土)】
- 【#CSHARP】System.Numerics.Vectorでグリングリン,AVXレジスタをぶん回したい 【2023年01月14日(土)】
- ビデオカードの大掃除をやってみた 【2012年02月09日(木)】
- Milkeyway@HomeのCUDA弾 【2010年12月14日(火)】
- グラボが届いた 【2010年10月14日(木)】
GPU上のグローバルメモリに,画像を流し込んで
- バイナリイメージにて
性能が出ない件(核爆
っちゅうか,CPUより遅いって何事じゃ,ぐるぅぁ~~ 😈
グローバルメモリとのやり取りには1往復600ナノ秒かかるとか何とか
これ,激しくアクセスすると,性能がた落ちは目に見えており...
共有メモリ16キロバイトを使おうかねぇ~
アルゴリズムが凄いことなるなぁ...
確実にバグるよなぁ...だったんすが
テクスチャメモリなるものをハケーン! 😎
こやつを使うと,キャッシュが効くおかげで,倍速以上は出るらしい...
っちゅうことで,32ビット画像にてテクスチャメモリを使ってみた
サンプルソースはこんな感じ?
m_ImageBufに32ビット画像(BGRA順)がバイナリイメージにて格納されておりまする
texture<signed int, 2, cudaReadModeElementType> g_tex; //テクスチャ変数の定義 cudaArray *g_cu_ary; cudaChannelFormatDesc g_channelDesc; //--------- void func() { g_channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindSigned); //テクスチャサイズの定義 CUDA_SAFE_CALL(cudaMallocArray(&g_cu_ary, &g_channelDesc, m_Width, m_Height)); //テクスチャ用の領域確保 CUDA_SAFE_CALL(cudaMemcpyToArray(g_cu_ary, 0, 0, m_ImageBuf, m_ImageBufLen, cudaMemcpyHostToDevice)); //テクスチャ領域に画像を転送 CUDA_SAFE_CALL(cudaBindTextureToArray(g_tex, g_cu_ary, g_channelDesc)); //テクスチャにバインド dim3 threads(16, 1, 1); dim3 grid(16,1,1); JobCoreTex<<<grid, threads>>>(); //カーネルの呼び出し CUDA_SAFE_CALL(cudaFreeArray(g_cu_ary)); //テクスチャ向け領域を破棄 }カーネルの中で,ピクセル値を取り出すときは
int px = tex2D(g_tex, x, y);
にてOK!
これで,アホみたいに速くなったぞぉ~
あと,カーネル内で,構造体ポインタからのメンバ参照とか,繁盛にやっちゃだめっす
特にループの中で,メンバ参照とかやったら,確実に泣きます
for(int i = 0; i < 100; i++) { int datX = tmpRec->point.X; int datY = tmpRec->point.Y; int px = tex2D(g_tex, x + i, y); }とか?
メイッパイ性能が落ちまする
CUDAコンパイラの最適化が甘い気がしますです
構造体ポインタのメンバはループの外で,
通常の変数に代入の上で,ループ突入したほうが幸せになれると思いますわん
っちゅうか,クラスやら構造体ポインタのメンバ参照に使う矢印(これ->ね)
正式名称,アロー演算子って言うっちゃ 😯
27年コンピュータ触ってきて,今日初めて知った(隕石爆
コメント
急激に寒くなってキター 晩秋がやってきた!
【#諸塚村】連休初日のBBQの模様なぞをアップしてみる
冷房なしで日中過ごすの,いつ以来よっ!?優しい涼しさが心地よい
【#異星人 発見!】異星人から謎の暗号メールを受け取った!解読不能ヤバス【シレペヒキク】
タッパーウェアが破産申請を計画、事業てこ入れ策つまずく-関係者
電子レンジが不調になったヤバス
【#地震】霧島山で有感地震 生駒高原のあたりか?(24/9/6)
奄美大島のマングース 環境省が根絶宣言