CUDAイメージローテーション

Question

CUDAでイメージローテーションを実装する際に問題があります。次のように私は非常に単純なローテート機能作業があります。

__device__ float readPixVal(float* ImgSrc,int ImgWidth,int x,int y) 
{ 
    return (float)ImgSrc[y*ImgWidth+x]; 
} 
__device__ void putPixVal(float* ImgSrc,int ImgWidth,int x,int y, float floatVal) 
{ 
    ImgSrc[y*ImgWidth+x] = floatVal; 
} 

__global__ void Rotate(float* Source, float* Destination, int sizeX, int sizeY, float deg) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x;// Kernel definition 
    int j = blockIdx.y * blockDim.y + threadIdx.y; 

    if(i < sizeX && j < sizeY) 
    { 
     putPixVal(Destination, sizeX, ((float)i)*cos(deg) - ((float)j)*sin(deg), ((float)i)*sin(deg) + ((float)j)*cos(deg)), readPixVal(Source, sizeX, i, j)); 
    } 
} 

問題は、私はすべての補間を行う方法がわからない、です。上記では、整数の丸めによって多くのピクセルがスキップされます。誰でもこの問題を解決する方法を知っているか、イメージのフリー/オープンソースの実装がありますか？私はCUDAを見つけることができませんでした。

Answer 1

これはトリック

__global__ void Rotate(float* Source, float* Destination, int sizeX, int sizeY, float deg) 
{ 
    int i = blockIdx.x * blockDim.x + threadIdx.x;// Kernel definition 
    int j = blockIdx.y * blockDim.y + threadIdx.y; 
    int xc = sizeX - sizeX/2; 
    int yc = sizeY - sizeY/2; 
    int newx = ((float)i-xc)*cos(deg) - ((float)j-yc)*sin(deg) + xc; 
    int newy = ((float)i-xc)*sin(deg) + ((float)j-yc)*cos(deg) + yc; 
    if (newx >= 0 && newx < sizeX && newy >= 0 && newy < sizeY) 
    { 
     putPixVal(Destination, sizeX, i , j, readPixVal(Source, sizeX, newx, newy)); 
    } 
} 

Answer 2

一般的に、この種の画像操作では、ソース画像の対応するピクセル（またはピクセルの補間グループ）を計算するすべての宛先ピクセル位置をループします。

これにより、通常あなたが気にする結果のイメージを均一かつ均一に塗りつぶすことができます。

Answer 3

void rotateImage_Kernel(cufftComplex* trg, const cufftComplex* src, const unsigned int imageWidth,const unsigned int imageHeight, const float angle, const float scale) 
{ 
    // compute thread dimension 
    const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
    const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 

    //// compute target address 
    const unsigned int idx = x + y * imageWidth; 

    const int xA = (x - imageWidth/2); 
    const int yA = (y - imageHeight/2); 

    const int xR = (int)floor(1.0f/scale * (xA * cos(angle) - yA * sin(angle))); 
    const int yR = (int)floor(1.0f/scale * (xA * sin(angle) + yA * cos(angle))); 

    float src_x = xR + imageWidth/2; 
    float src_y = yR + imageHeight/2; 



    if (src_x >= 0.0f && src_x < imageWidth && src_y >= 0.0f && src_y < imageHeight) { 
     // BI - LINEAR INTERPOLATION 
     float src_x0 = (float)(int)(src_x); 
     float src_x1 = (src_x0+1); 
     float src_y0 = (float)(int)(src_y); 
     float src_y1 = (src_y0+1); 

     float sx = (src_x-src_x0); 
     float sy = (src_y-src_y0); 


     int idx_src00 = min(max(0.0f,src_x0 + src_y0 * imageWidth),imageWidth*imageHeight-1.0f); 
     int idx_src10 = min(max(0.0f,src_x1 + src_y0 * imageWidth),imageWidth*imageHeight-1.0f); 
     int idx_src01 = min(max(0.0f,src_x0 + src_y1 * imageWidth),imageWidth*imageHeight-1.0f); 
     int idx_src11 = min(max(0.0f,src_x1 + src_y1 * imageWidth),imageWidth*imageHeight-1.0f); 

     trg[idx].y = 0.0f; 

     trg[idx].x = (1.0f-sx)*(1.0f-sy)*src[idx_src00].x; 
     trg[idx].x += ( sx)*(1.0f-sy)*src[idx_src10].x; 
     trg[idx].x += (1.0f-sx)*( sy)*src[idx_src01].x; 
     trg[idx].x += ( sx)*( sy)*src[idx_src11].x; 
    } else { 
     trg[idx].x = 0.0f; 
     trg[idx].y = 0.0f; 
    } 

    DEVICE_METHODE_LAST_COMMAND; 

} 


void translateImage_Kernel(cufftComplex* trg, const cufftComplex* src, const unsigned int imageWidth, const unsigned int imageHeight, const float tX, const float tY) 
{ 
    // compute thread dimension 
    const unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; 
    const unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; 

    //// compute target address 
    const unsigned int idx = x + y * imageWidth; 

    const int xB = ((int)x + (int)tX); 
    const int yB = ((int)y + (int)tY); 

    if (xB >= 0 && xB < imageWidth && yB >= 0 && yB < imageHeight) { 
     trg[idx] = src[xB + yB * imageWidth]; 
    } else { 
     trg[idx].x = 0.0f; 
     trg[idx].y = 0.0f; 
    } 

    DEVICE_METHODE_LAST_COMMAND; 

}