GpuMat - 사용자 정의 커널에서 2 채널 플로트 데이터에 액세스

Question

cv::cuda::PtrStep은 데이터를 사용자 정의 커널에 직접 전달하는 데 사용됩니다. 나는 하나의 채널 액세스 here 그러나 나의 경우 2 채널 매트 (CV_32FC2)의 예를 발견했다. 이 경우 복잡한 값이 다음과 같이 인코딩되는 복잡한 절대 제곱 값을 얻으려고합니다. 실수 부분은 첫 번째 평면이고, 허수 부분은 주어진 두 번째 평면입니다. Mat.

__global__ void testKernel(const cv::cuda::PtrStepSz<cv::Vec2f> input, cv::cuda::PtrStepf output) 
{ 
    int x = blockIdx.x * blockDim.x + threadIdx.x; 
    int y = blockIdx.y * blockDim.y + threadIdx.y; 

    if (x <= input.cols - 1 && y <= input.rows - 1 && y >= 0 && x >= 0) 
    { 
     float val_re = input(x, y)[0]; 
     float val_im = input(x, y) [1]; 
     output(x, y) = val_re * val_re + val_im * val_im; 
    } 
} 

을하지만, 이것은 다음과 같은 오류가 발생합니다 :

나는 시도 나는 그것을 얻을

calling a __host__ function("cv::Vec<float, (int)2> ::operator []") from a __global__ function("gpuholo::testKernel") is not allowed 

. []은 cv::Vec2f이 아니기 때문에 cv::cuda::Vec2f (분명히 존재하지 않음)이므로 __host__ 제한된 기능입니다. 하지만 여전히 데이터에 액세스하고 싶습니다.

Vec2f과 비슷한 장치 측의 2 채널 데이터에 액세스하는 다른 메커니즘이 있습니까? 커널의 모습 있도록

---

나는 두 CV_32FC1Mat들에 input 분할의 형태로 대안의 생각 :

__global__ void testKernel(const cv::cuda::PtrStepSzf re, const cv::cuda::PtrStepSzf im, cv::cuda::PtrStepf output) 

하지만 난 "청소기"솔루션은 Vec2f이 있는지 여부 궁금하네요 - 같은 것.

Answer 1

사용자 정의 CUDA 커널에서 원시 데이터 유형을 사용하여 GpuMat의 데이터에 액세스 할 수 있습니다. 예 : CUDA 런타임에서 제공하는 float2 유형은 cv::Vec2f의 부분 대체품으로 사용할 수 있습니다. 다음은 GpuMat 데이터에 액세스하기위한 원시 데이터 형식의 사용을 보여주는 예제 코드입니다.

#include <iostream> 
#include <cuda_runtime.h> 
#include <opencv2/opencv.hpp> 

using std::cout; 
using std::endl; 

__global__ void kernel_absolute(float2* src, float* dst, int rows, int cols, int iStep, int oStep) 
{ 
    int i = blockIdx.y * blockDim.y + threadIdx.y; //Row number 
    int j = blockIdx.x * blockDim.x + threadIdx.x; //Column number 

    if (i<rows && j<cols) 
    { 
     /* Compute linear index from 2D indices */ 
     int tidIn = i * iStep + j; 
     int tidOut = i * oStep + j; 

     /* Read input value */ 
     float2 input = src[tidIn]; 

     /* Calculate absolute value */ 
     float output = sqrtf(input.x * input.x + input.y * input.y); 

     /* Write output value */ 
     dst[tidOut] = output; 
    } 
} 

int main(int argc, char** argv) 
{ 
    /* Example to compute absolute value of each element of a complex matrix */ 
    int rows = 10; 
    int cols = 10; 
    int input_data_type = CV_32FC2; //input is complex 
    int output_data_type = CV_32FC1; //output is real 

    /* Create input matrix on host */ 
    cv::Mat input = cv::Mat::zeros(rows,cols,input_data_type) + cv::Vec2f(1,1) /* Initial value is (1,1) */; 

    /* Display input */ 
    cout<<input<<endl; 

    /* Create input matrix on device */ 
    cv::cuda::GpuMat input_d; 
    /* Copy from host to device */ 
    input_d.upload(input); 

    /* Create output matrix on device */ 
    cv::cuda::GpuMat output_d(rows,cols, output_data_type); 

    /* Compute element step value of input and output */ 
    int iStep = input_d.step/sizeof(float2); 
    int oStep = output_d.step/sizeof(float); 

    /* Choose appropriate block size */ 
    dim3 block(8,8); 

    /* Compute grid size using input size and block size */ 
    dim3 grid ((cols + block.x -1)/block.x, (rows + block.y -1)/block.y); 

    /* Launch CUDA kernel to compute absolute value */ 
    kernel_absolute<<<grid, block>>>(reinterpret_cast<float2*>(input_d.data), reinterpret_cast<float*>(output_d.data), rows, cols, iStep, oStep); 

    /* Check kernel launch errors */ 
    assert(cudaSuccess == cudaDeviceSynchronize()); 

    cv::Mat output; 

    /* Copy results from device to host */ 
    output_d.download(output); 

    /* Display output */ 
    cout<<endl<<output<<endl; 

    return 0; 
} 

집계하고 CUDA 8.0 우분투 14.04에서 다음 명령을 사용하여 테스트 :

nvcc -o complex complex.cu -arch=sm_61 -L/usr/local/lib -lopencv_core

Answer 2

커널에 단일 입력으로 작업하려면 flatten your 2 channel image to a 1 channel image을 사용할 수 있습니다.

// test image 
Mat h_mat(Size(50,50),CV_32FC2,Scalar(0.0)); 

// Mat::reshape takes number of channels and rows, for your example 1,1 
Mat h_mat_flat = h_mat.reshape(1,1); 

// to upload to gpu 
GpuMat d_mat_flat(h_mat_flat.size(), h_mat_flat.type()); 
d_mat_flat.upload(h_mat_flat); 

이제 PtrStepSzf으로 커널에 d_mat_flat를 전달할 수 있습니다.