バイリニア補間を使用して画像のサイズを変更するためのCUDA関数を実装しました。出力画像の正確な値を確認するために小さな行列でテストするまで、この関数はおそらく正しい結果を(視覚的に)提供していました。私が得た結果は、OpenCVやMATLABの結果とは異なりました。アルゴリズムに明らかな欠陥が見つかりません。誰かがこれについて私を助けることができますか?
バイリニア補間デバイス機能:
texture<float, cudaTextureType2D> tex32f;
//Device function
__device__ float blinterp(const float xIndex, const float yIndex)
{
//floor the coordinates to get to the nearest valid pixel
const int intX = static_cast<int>(xIndex);
const int intY = static_cast<int>(yIndex);
//Set weights of pixels according to distance from actual location
const float a = xIndex - intX;
const float b = yIndex - intY;
/* _____________________
*| | |
*|(1-a)(1-b)| (a)(1-b) |
*|__________|__________|
*| | |
*| (1-a)(b) | (a)(b) |
*|__________|__________|
*/
//Compute the weighted average of 4 nearest pixels
float out = (1 - a) * (1 - b) * tex2D(tex32f, intX,intY)
+ (a) * (1 - b) * tex2D(tex32f,intX + 1,intY)
+ (1 - a) * (b) * tex2D(tex32f, intX,intY + 1)
+ (a * b) * tex2D(tex32f,intX + 1,intY + 1);
return out;
}
カーネルのサイズ変更:
__global__ void kernel_resize(float* dst, int dstWidth, int dstHeight, int dstPitch, float xScale, float yScale)
{
const int xIndex = blockIdx.x * blockDim.x + threadIdx.x;
const int yIndex = blockIdx.y * blockDim.y + threadIdx.y;
if(xIndex>=dstWidth || yIndex>=dstHeight) return;
const unsigned int tid = yIndex * dstPitch + xIndex;
const float inXindex = xIndex * xScale;
const float inYindex = yIndex * yScale;
dst[tid] = blinterp(inXindex,inYindex);
}
ラッパー関数:
int resize_32f_c1(float* src,float* dst,int srcWidth,int srcHeight, int srcPitch, int dstWidth,int dstHeight,int dstPitch)
{
if((srcWidth == dstWidth) && (srcHeight == dstHeight))
{
cudaMemcpy2D(dst,dstPitch,src,srcPitch,srcWidth * sizeof(float),srcHeight,cudaMemcpyDeviceToDevice);
return 0;
}
cudaBindTexture2D(NULL,tex32f,src,srcWidth,srcHeight,srcPitch);
dim3 Block(16,16);
dim3 Grid((dstWidth + Block.x - 1)/Block.x, (dstHeight + Block.y - 1)/Block.y);
float x = (float)(srcWidth)/(float)dstWidth;
float y = (float)(srcHeight)/(float)dstHeight;
kernel_resize<<<Grid,Block>>>(dst,dstWidth,dstHeight,dstPitch/sizeof(float),x,y);
cudaUnbindTexture(tex32f);
return 0;
}
結果(2ダウンスケール):
入力(10 x 10):
0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0
0 0 0 1 1 1 1 0 0 0
0 0 0 1 1 1 1 0 0 0
0 0 0 1 1 1 1 0 0 0
0 0 0 1 1 1 1 0 0 0
0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0
0 0 0 0 0 0 0 0 0 0
MATLABおよびOpenCV出力:
0 0 0 0 0
0 0.25 0.5 0.25 0
0 0.5 1 0.5 0
0 0.25 0.5 0.25 0
0 0 0 0 0
私の出力:
0 0 0 0 0
0 0 0 0 0
0 0 1 1 0
0 0 1 1 0
0 0 0 0 0