3

私は 3x3 ソーベル フィルター用の OpenCL カーネルを作成しました。現在、2k x 2k のイメージで約 17 ミリ秒で実行されています。これは私が望んでいたほどの断食ではありません。速度を改善する方法について何か提案はありますか? カーネルを最適化するためのチェックリストの提案のほとんどに従いました。私のプロセッサは Intel i5-3450 です。ワークグループのサイズは 8x8 で、ワークアイテムの数は高さ x 幅 / 16 で、実行中の画像では 2048 x 128 です。

__kernel void localCacheSobelFilter(
  const __global char16* src, 
  __write_only __global float16* angle, 
  __write_only __global float16* mag, 
  const int width, 
  const int height)
{
  // Cache the data we're looking at in __local space
  const int row = get_global_id(0);
  const int col = get_global_id(1);

  const int cacheRow = get_local_id(0) + 1;
  const int cacheCol = get_local_id(1) + 1;

  __local char16 cache[BLOCK_SIZE + 2][BLOCK_SIZE + 2];


  cache[cacheRow][cacheCol] = src[ indexOf(row, col) ];

  // --- Deal with the boundary conditions
  // This adds in the rows above and below the local block,
  // ignoring the corners.
  const bool atTopRow = (cacheRow == 1);
  const bool atBottomRow = (cacheRow == BLOCK_SIZE);

  if(atTopRow) {
    cache[0][cacheCol] = src[ indexOf(row - 1, col) ];

  } else if (atBottomRow) {
    cache[BLOCK_SIZE + 1][cacheCol] = src[ indexOf(row + 1, col) ];
  }


  // This adds in the columns to the left and right of the local block,
  // ignoring the corners.
  const bool atLeftCol = (cacheCol == 1);
  const bool atRightCol = (cacheCol == BLOCK_SIZE);

  if(atLeftCol) { 
    cache[cacheRow][0].sf = src[ indexOf(row, col - 1) ].sf;

  } else if (atRightCol) {
    cache[cacheRow][BLOCK_SIZE + 1].s0 = src[ indexOf(row, col + 1) ].s0;
  }


  // Now finally check the corners
  const bool atTLCorner = atTopRow && atLeftCol;
  const bool atTRCorner = atTopRow && atRightCol;
  const bool atBLCorner = atBottomRow && atLeftCol;
  const bool atBRCorner = atBottomRow && atRightCol;


  if(atTLCorner) {
    cache[0][0].sf = src[ indexOf(row - 1, col - 1) ].sf;

  } else if (atTRCorner) { 
    cache[0][BLOCK_SIZE + 1].s0 = src[ indexOf(row - 1, col + 1) ].s0;

  } else if (atBLCorner) {
    cache[BLOCK_SIZE + 1][0].sf = src[ indexOf(row + 1, col - 1) ].sf;

  } else if (atBRCorner) {
    cache[BLOCK_SIZE + 1][BLOCK_SIZE + 1].s0 = src[ indexOf(row + 1, col + 1) ].s0;
  }

  barrier(CLK_LOCAL_MEM_FENCE); 

  //===========================================================================
  // Do the calculation

  //  [..., pix00]  upperRow  [pix02, ...]
  //  [..., pix10]  centerRow [pix12, ...]
  //  [..., pix20]  lowerRow  [pix22, ...]
  const char pix00 = cache[cacheRow - 1][cacheCol - 1].sf;
  const char pix10 = cache[cacheRow    ][cacheCol - 1].sf;
  const char pix20 = cache[cacheRow + 1][cacheCol - 1].sf;

  const char16 upperRow  = cache[cacheRow - 1][cacheCol];    
  const char16 centerRow = cache[cacheRow    ][cacheCol];
  const char16 lowerRow  = cache[cacheRow + 1][cacheCol];

  const char pix02 = cache[cacheRow - 1][cacheCol + 1].s0;
  const char pix12 = cache[cacheRow    ][cacheCol + 1].s0;              
  const char pix22 = cache[cacheRow + 1][cacheCol + 1].s0;



  // Do the calculations for Gy
  const char16 upperRowShiftLeft = (char16)(upperRow.s123456789abcdef, pix02);
  const char16 upperRowShiftRight = (char16)(pix00, upperRow.s0123456789abcde);

  const char16 lowerRowShiftLeft = (char16)(lowerRow.s123456789abcdef, pix22);
  const char16 lowerRowShiftRight = (char16)(pix20, lowerRow.s0123456789abcde);

  const float16 Gy = convert_float16(
    (upperRowShiftLeft + 2 * upperRow + upperRowShiftRight)
    - (lowerRowShiftLeft + 2 * lowerRow + lowerRowShiftRight));


  // Do the calculations for Gx
  const char16 centerRowShiftLeft = (char16)(centerRow.s123456789abcdef, pix12);
  const char16 centerRowShiftRight = (char16)(pix10, centerRow.s0123456789abcde);

  const float16 Gx = convert_float16(
    (upperRowShiftRight + 2 * centerRowShiftRight + lowerRowShiftRight)
    - (upperRowShiftLeft + 2 * centerRowShiftLeft + lowerRowShiftLeft));


  // Find the angle and magnitude
  angle[ indexOf(row, col) ] = 0.0; //atan2(Gy, Gx);
  mag[ indexOf(row, col) ] = ALPHA * max(Gx, Gy) + BETA * min(Gx, Gy);
}

どんな助けでも大歓迎です。ありがとう！

4

1 に答える 1

2

カーネルには、エッジ効果を回避するための「if」がたくさんあります。しかし、99％の時間は時間がかかり、役に立たない。NDRangeKernelでグローバル作業サイズを変更し、オフセットを使用してエッジ効果を回避できると思います。

例：オフセット= {1,1,0} GlobalWorkSize = {width-2、height-2,0}

于 2012-10-16T07:38:20.350 に答える