Вычитание изображения с помощью CUDA и текстур

Question

Моя цель - использовать C ++ с CUDA, чтобы вычесть темную рамку из необработанного изображения. Я хочу использовать текстуры для ускорения. Ввод изображений - cv :: Mat с типом CV_8UC4 (я использую указатель на данные cv :: Mat). Это ядро, которое я придумал, но я не знаю, как со временем вычесть текстуры друг из друга:

__global__ void DarkFrameSubtractionKernel(unsigned char* outputImage, size_t pitchOutputImage,
cudaTextureObject_t inputImage, cudaTextureObject_t darkImage, int width, int height)
{
    const int x = blockIdx.x * blockDim.x + threadIdx.x;
    const int y = blockDim.y * blockIdx.y + threadIdx.y;

    const float tx = (x + 0.5f);
    const float ty = (y + 0.5f);

    if (x >= width || y >= height) return;

    uchar4 inputImageTemp = tex2D<uchar4>(inputImage, tx, ty);
    uchar4 darkImageTemp = tex2D<uchar4>(darkImage, tx, ty);

    outputImage[y * pitchOutputImage + x] = inputImageTemp - darkImageTemp; // this line will throw an error
}

Это функция, которая вызывает ядро ​​(вы можете видеть, что я создаю текстуры from unsigned char):

void subtractDarkImage(unsigned char* inputImage, size_t pitchInputImage, unsigned char* outputImage,
size_t pitchOutputImage, unsigned char* darkImage, size_t pitchDarkImage, int width, int height, 
cudaStream_t stream)
{
    cudaResourceDesc resDesc = {};
    resDesc.resType = cudaResourceTypePitch2D;
    resDesc.res.pitch2D.width = width;
    resDesc.res.pitch2D.height = height;
    resDesc.res.pitch2D.devPtr = inputImage;
    resDesc.res.pitch2D.pitchInBytes = pitchInputImage;
    resDesc.res.pitch2D.desc = cudaCreateChannelDesc(8, 8, 8, 8, cudaChannelFormatKindUnsigned);

    cudaTextureDesc texDesc = {};
    texDesc.readMode = cudaReadModeElementType;
    texDesc.addressMode[0] = cudaAddressModeBorder;
    texDesc.addressMode[1] = cudaAddressModeBorder;

    cudaTextureObject_t imageInputTex, imageDarkTex;
    CUDA_CHECK(cudaCreateTextureObject(&imageInputTex, &resDesc, &texDesc, 0));

    resDesc.res.pitch2D.devPtr = darkImage;
    resDesc.res.pitch2D.pitchInBytes = pitchDarkImage;
    CUDA_CHECK(cudaCreateTextureObject(&imageDarkTex, &resDesc, &texDesc, 0));

    dim3 block(32, 8);
    dim3 grid = paddedGrid(block.x, block.y, width, height);

    DarkImageSubtractionKernel << <grid, block, 0, stream >> > (reinterpret_cast<uchar4*>(outputImage), pitchOutputImage / sizeof(uchar4),
    imageInputTex, imageDarkTex, width, height);


    CUDA_CHECK(cudaDestroyTextureObject(imageInputTex));
    CUDA_CHECK(cudaDestroyTextureObject(imageDarkTex));
}

Код не компилируется, так как я не могу вычесть uchar4 из другого (в ядре). Есть ли простой способ вычитания здесь?

Помощь очень ценится.

Answer 1

Есть ли здесь простой способ вычитания?

Не существует арифметических c операторов, определенных для встроенных векторных типов CUDA. Если вы замените

outputImage[y * pitchOutputImage + x] = inputImageTemp - darkImageTemp;

на

uchar4 val;
val.x = inputImageTemp.x - darkImageTemp.x;
val.y = inputImageTemp.y - darkImageTemp.y;
val.z = inputImageTemp.z - darkImageTemp.z;
val.w = inputImageTemp.w - darkImageTemp.w;
outputImage[y * pitchOutputImage + x] = val;

, все будет работать. Если это вас оскорбляет, я предлагаю написать небольшую библиотеку вспомогательных функций, чтобы скрыть беспорядок.