플로트가있는 사용자 정의 커널 GpuMat

Question

GpuMat 데이터를 사용하여 사용자 정의 커널을 작성하여 이미지의 픽셀의 아크 코사인을 찾으려고합니다. GPU에 CV_8UC1 데이터가 있지만 데이터를 업로드 할 때 값을 업로드, 다운로드 및 변경할 수는 있지만 문자는 아크 코사인을 계산하는 데 사용할 수 없습니다. 그러나 내 GPU를 CV_32FC1 유형 (부동 소수점)으로 변환하려고하면 다운로드 부분에서 잘못된 메모리 액세스 오류가 발생합니다. 여기 내 코드입니다 :

//.cu code 
#include <cuda_runtime.h> 
#include <stdlib.h> 
#include <iostream> 
#include <stdio.h> 
__global__ void funcKernel(const float* srcptr, float* dstptr, size_t srcstep, const  size_t dststep, int cols, int rows){ 
    int rowInd = blockIdx.y*blockDim.y+threadIdx.y; 
    int colInd = blockIdx.x*blockDim.x+threadIdx.x; 
    if(rowInd >= rows || colInd >= cols) 
      return; 
    const float* rowsrcptr=srcptr+rowInd*srcstep; 
    float* rowdstPtr= dstptr+rowInd*dststep; 
    float val = rowsrcptr[colInd]; 
    if((int) val % 90 == 0) 
      rowdstPtr[colInd] = -1 ; 
    else{ 
      float acos_val = acos(val); 
      rowdstPtr[colInd] = acos_val; 
    } 
} 

int divUp(int a, int b){ 
    return (a+b-1)/b; 
} 

extern "C" 
{ 
void func(const float* srcptr, float* dstptr, size_t srcstep, const size_t dststep, int cols, int rows){ 
    dim3 blDim(32,8); 
    dim3 grDim(divUp(cols, blDim.x), divUp(rows,blDim.y)); 
    std::cout << "calling kernel from func\n"; 
    funcKernel<<<grDim,blDim>>>(srcptr,dstptr,srcstep,dststep,cols,rows); 
    std::cout << "done with kernel call\n"; 
    cudaDeviceSynchronize(); 
} 

//.cpp code 
void callKernel(const GpuMat &src, GpuMat &dst){ 
    float* p = (float*)src.data; 
    float* p2 =(float*) dst.data; 
    func(p,p2,src.step,dst.step,src.cols,src.rows); 
} 

int main(){ 
    Mat input = imread("cat.jpg",0); 
    Mat float_input; 
    input.convertTo(float_input,CV_32FC1); 
    GpuMat d_frame,d_output; 
    Size size = float_input.size(); 
    d_frame.upload(float_input); 
    d_output.create(size,CV_32FC1); 
    callKernel(d_frame,d_output); 
    Mat output(d_output); 
    return 0; 
} 

내가 프로그램을 실행하면 내 컴파일러는 나에게 알려줍니다이 :

OpenCV Error: Gpu API call (an illegal memory access was encountered) in copy, file /home/mobile/opencv-2.4.9/modules/dynamicuda/include/opencv2/dynamicuda/dynamicuda.hpp, line 882 terminate called after throwing an instance of 'cv::Exception' what(): /home/mobile/opencv-2.4.9/modules/dynamicuda/include/opencv2/dynamicuda/dynamicuda.hpp:882: error: (-217) an illegal memory access was encountered in function copy

Answer 1

는 오프셋 float 인 것처럼 당신은 이미지 step을 치료하고 있습니다. 한 행에서 다음 행까지의 바이트 오프셋입니다.

이 대신 같은 것을보십시오 :

const float* rowsrcptr= (const float *)(((char *)srcptr)+rowInd*srcstep); 
float* rowdstPtr= (float *) (((char *)dstptr)+rowInd*dststep); 

을 documentation에서 :

step – Number of bytes each matrix row occupies.

또한 (예를 들어, func에) 코드에 proper cuda error checking를 추가하는 것이 좋습니다. cuda-memcheck으로 코드를 실행하면 실제 커널 오류로 인해 잘못된 읽기/쓰기가 발생합니다.

Answer 2

당신은 당신의 자신의 커널을 작성하는 cv::cuda::PtrStp<> 또는 cv::cuda::PtrStpSz<>을 사용할 수 있습니다 (그래서 당신이 GpuMat의 스텝 매개 변수를 사용하지 않은 당신의 코드 litte 조금 단순화 : D) :

커널 :

을

__global__ void myKernel(const cv::cuda::PtrStepSzf input, 
          cv::cuda::PtrStepSzf 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) 
     { 
      output(y, x) = input(y, x); 
     } 
    } 

주의 사항 :
cv::cuda::PtrStep<> :
cv::cuda::PtrStepSz<> 크기 정보가없는 : 크기 정보와
cv::cuda::PtrStepSzb : 서명 숯불 엄마를위한 TS (CV_8U)
cv::cuda::PtrStepSzf : 다른 유형의 예를

커널 호출 :

void callKernel(cv::InputArray _input, 
        cv::OutputArray _output, 
        cv::cuda::Stream _stream) 
    { 
     const cv::cuda::GpuMat input = _input.getGpuMat(); 

     _output.create(input.size(), input.type()); 
     cv::cuda::GpuMat output = _output.getGpuMat(); 

     dim3 cthreads(16, 16); 
     dim3 cblocks(
      static_cast<int>(std::ceil(input1.size().width/
       static_cast<double>(cthreads.x))), 
      static_cast<int>(std::ceil(input1.size().height/
       static_cast<double>(cthreads.y)))); 

     cudaStream_t stream = cv::cuda::StreamAccessor::getStream(_stream); 
     myKernel<<<cblocks, cthreads, 0, stream>>>(input, output); 

     cudaSafeCall(cudaGetLastError()); 
    } 

당신은 cv::cuda::GpuMat을 사용하여이 funktion를 호출 할 수

callKernel(d_frame, d_output, cv::cuda::Stream()); 

플로트 매트 (CV_32F)
cv::cuda::PtrStep<cv::Point2f>에 대한