CUDAの画像処理はテクスチャが必須ですな

ＧＰＵ上のグローバルメモリに，画像を流し込んで

バイナリイメージにて

しこしこ，画像解析をやっておりましたが．．．

性能が出ない件(核爆

っちゅうか，ＣＰＵより遅いって何事じゃ，ぐるぅぁ～～ 😈


グローバルメモリとのやり取りには１往復600ナノ秒かかるとか何とか
これ，激しくアクセスすると，性能がた落ちは目に見えており．．．

共有メモリ１６キロバイトを使おうかねぇ～
アルゴリズムが凄いことなるなぁ．．．
確実にバグるよなぁ．．．だったんすが

テクスチャメモリなるものをハケーン！ 😎
こやつを使うと，キャッシュが効くおかげで，倍速以上は出るらしい．．．

っちゅうことで，３２ビット画像にてテクスチャメモリを使ってみた
サンプルソースはこんな感じ？

m_ImageBufに３２ビット画像(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);
にてＯＫ！
これで，アホみたいに速くなったぞぉ～

あと，カーネル内で，構造体ポインタからのメンバ参照とか，繁盛にやっちゃだめっす
特にﾙｰﾌﾟの中で，メンバ参照とかやったら，確実に泣きます

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コンパイラの最適化が甘い気がしますです

構造体ポインタのメンバはﾙｰﾌﾟの外で，
通常の変数に代入の上で，ﾙｰﾌﾟ突入したほうが幸せになれると思いますわん

っちゅうか，クラスやら構造体ポインタのメンバ参照に使う矢印(これ->ね）
正式名称，アロー演算子って言うっちゃ 😯


２７年コンピュータ触ってきて，今日初めて知った(隕石爆