cuda . के साथ bgr को yuv420p में बदलने की समस्या

विभिन्न प्रकार के मुद्दे हैं:

• आपके सीपीयू और जीपीयू कोड के बीच आर, जी, बी को वाई, यू, वी में बदलने की गणना समान नहीं है। हाँ, यह मायने रखता है।
• आपके सीपीयू कोड में प्लानर वाई, यू, वी स्टोरेज है। इसका मतलब है कि वाई का अपना विमान है, यू का अपना विमान है, और वी का अपना विमान है। आपका GPU कोड सेमी प्लानर (NV12) प्रारूप है। इसका मतलब है कि वाई का अपना विमान है, और यू, वी एक ही विमान में अंतःस्थापित हैं: यूवीयूवीयूवीयूवीयूवीयूवी .... जाहिर है कि उन दो कोडों का आउटपुट कभी भी समान रूप से मेल नहीं खा सकता है।
• IMO, इसमें OpenCV को खींचने की कोई आवश्यकता नहीं है।
• कर्नेल (GPU) कोड में आपकी UV ऑफ़सेट गणना टूट गई थी। imgwidth*imgheight ऑफसेट आपको Y क्षेत्र (सही ढंग से) से आगे ले जाता है, लेकिन उस बिंदु से, यूवी प्लानर क्षेत्र में पंक्ति द्वारा अनुक्रमित करने के लिए row_num*imgwidth का उपयोग करना सही नहीं है। आपके पास यूवी प्लानर क्षेत्र में इतनी पंक्तियाँ नहीं हैं, आपके पास केवल आधी पंक्तियाँ हैं।
• आपके GPU कर्नेल में, आपके पास U,V ऑर्डर उल्टा था, आप प्रभावी रूप से VUVUVUVU कर रहे थे ...

मेरी सिफारिश गणना अंतर और भंडारण आदेश/प्रारूप के सामंजस्य से शुरू करने की होगी। निम्नलिखित कोड में उपरोक्त मुद्दों को संबोधित किया गया है, और सीपीयू और जीपीयू कोड के बीच मेरे लिए मिलान परिणाम देता है:

$ cat t1708.cu
#include <iostream>
#include <time.h>
#include <cstdlib>
using namespace std;
// I have no idea if these are the correct conversion formulas
// I simply lifted what I saw in your host code so that we 
// are using the same conversion calculations in host and device
__host__ __device__ unsigned char bgr2y(int R, int G, int B){
  int Y = ((66 * R + 129 * G + 25 * B + 128) >> 8) + 16;
  return (unsigned char)((Y<0)? 0 : ((Y > 255) ? 255 : Y));}
__host__ __device__ int bgr2u(int R, int G, int B){
  int U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128;
  return (unsigned char)((U<0)? 0 : ((U > 255) ? 255 : U));}
__host__ __device__ int bgr2v(int R, int G, int B){
  int V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128;
  return (unsigned char)((V<0)? 0 : ((V > 255) ? 255 : V));}

void bgr_to_yuv420p(unsigned  char* yuv420p, unsigned char* bgr, int width, int height)
{
    if (yuv420p == NULL || bgr== NULL)
        return;
    int frameSize = width*height;

    int yIndex = 0;
    int uIndex = frameSize;

    int R, G, B;
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            B = bgr[(i * width + j) * 3 + 0];
            G = bgr[(i * width + j) * 3 + 1];
            R = bgr[(i * width + j) * 3 + 2];

            //BGR to YUV
            yuv420p[yIndex++] = bgr2y(R,G,B);
            if (i % 2 == 0 && j % 2 == 0)
            {
                yuv420p[uIndex] = bgr2u(R,G,B);
                yuv420p[uIndex+1] = bgr2v(R,G,B);
                uIndex+=2;
            }
        }
    }
}

//kernel function to convert bgr to yuv420p
__global__ void bgr2yuv420p(uchar3 *  d_in, unsigned char * d_out,
                               uint imgheight, uint imgwidth)
{

    int col_num = blockIdx.x*blockDim.x+threadIdx.x;
    int row_num = blockIdx.y*blockDim.y+threadIdx.y;

    if ((row_num < imgheight) && (col_num < imgwidth))
    {
//        uint32_t a = *((uint32_t *)&dinput[global_offset*3]);
        int global_offset = row_num*imgwidth+col_num;

        int r,g,b;
        r = int(d_in[global_offset].z);
        g = int (d_in[global_offset].y);
        b = int (d_in[global_offset].x);


        d_out[row_num * imgwidth + col_num] = bgr2y(r,g,b);
        if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
            int uv_offset = imgwidth*imgheight+((row_num>>1)*imgwidth)+col_num;
            d_out[uv_offset] = bgr2u(r,g,b);
            d_out[uv_offset+1] = bgr2v(r,g,b);

        }

    }
}

int main(void)
{

    const uint imgheight = 1000;
    const uint imgwidth = 1500;

    //input and output
    uchar3 *d_in;
    unsigned char *d_out;
    uchar3 *idata = new uchar3[imgheight*imgwidth];
    unsigned char *odata = new unsigned char[imgheight*imgwidth*3/2];
    unsigned char *cdata = new unsigned char[imgheight*imgwidth*3/2];
    uchar3 pix;
    for (int i = 0; i < imgheight*imgwidth; i++){
      pix.x = (rand()%30)+40;
      pix.y = (rand()%30)+40;
      pix.z = (rand()%30)+40;
      idata[i] = pix;}
    for (int i = 0; i < imgheight*imgwidth; i++) idata[i] = pix;
    bgr_to_yuv420p(cdata, (unsigned char*) idata, imgwidth, imgheight);
    // malloc memo in gpu
    cudaMalloc((void**)&d_in, imgheight*imgwidth*sizeof(uchar3));
    cudaMalloc((void**)&d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2);

    //copy image from cpu to gpu
    cudaMemcpy(d_in, idata, imgheight*imgwidth*sizeof(uchar3), cudaMemcpyHostToDevice);

    dim3 threadsPerBlock(32, 32);
    dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x,
                       (imgheight + threadsPerBlock.y - 1) / threadsPerBlock.y);

    //run kernel function
    bgr2yuv420p<<<blocksPerGrid, threadsPerBlock>>>(d_in, d_out, imgheight, imgwidth);

    cudaDeviceSynchronize();

    //copy yuv420p from gpu to cpu
    cudaMemcpy(odata, d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2, cudaMemcpyDeviceToHost);
    for (int i = 0; i < (imgwidth*imgheight*3/2); i++) if (odata[i] != cdata[i]) {std::cout << "mismatch at: " << i << " was: " << (int)odata[i] << " should be: " << (int)cdata[i] << std::endl; return 0;}
    cudaFree(d_in);
    cudaFree(d_out);


    return 0;

}
$ nvcc -o t1708 t1708.cu
$ cuda-memcheck ./t1708
========= CUDA-MEMCHECK
========= ERROR SUMMARY: 0 errors
$

किसी भी समय आपको CUDA कोड से परेशानी हो रही है, मैं अनुशंसा करता हूं

1. उचित CUDA त्रुटि जाँच
2. cuda-memcheck के साथ अपना कोड चलाना

संपादित करें: अतिरिक्त टिप्पणियों के आधार पर, यहां उपरोक्त कोड का एक संस्करण है जो ओपी द्वारा आपूर्ति किए गए सीपीयू कोड शब्दशः का उपयोग करता है, और एक CUDA कर्नेल प्रदान करता है जो YUV प्लानर स्टोरेज (सेमी-प्लानर स्टोरेज के बजाय) उत्पन्न करता है:

#include <iostream>
#include <time.h>
#include <cstdlib>
using namespace std;
__host__ __device__ unsigned char bgr2y(int R, int G, int B){
  int Y = ((66 * R + 129 * G + 25 * B + 128) >> 8) + 16;
  return (unsigned char)((Y<0)? 0 : ((Y > 255) ? 255 : Y));}
__host__ __device__ int bgr2u(int R, int G, int B){
  int U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128;
  return (unsigned char)((U<0)? 0 : ((U > 255) ? 255 : U));}
__host__ __device__ int bgr2v(int R, int G, int B){
  int V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128;
  return (unsigned char)((V<0)? 0 : ((V > 255) ? 255 : V));}

void bgr_to_yuv420sp(unsigned  char* yuv420p, unsigned char* bgr, int width, int height)
{
    if (yuv420p == NULL || bgr== NULL)
        return;
    int frameSize = width*height;

    int yIndex = 0;
    int uIndex = frameSize;

    int R, G, B;
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            B = bgr[(i * width + j) * 3 + 0];
            G = bgr[(i * width + j) * 3 + 1];
            R = bgr[(i * width + j) * 3 + 2];

            //BGR to YUV
            yuv420p[yIndex++] = bgr2y(R,G,B);
            if (i % 2 == 0 && j % 2 == 0)
            {
                yuv420p[uIndex] = bgr2u(R,G,B);
                yuv420p[uIndex+1] = bgr2v(R,G,B);
                uIndex+=2;
            }
        }
    }
}
void bgr_to_yuv420p(unsigned  char* yuv420p, unsigned char* bgr, int width, int height)
{
    if (yuv420p == NULL || bgr== NULL)
        return;
    int frameSize = width*height;
    int chromaSize = frameSize / 4;

    int yIndex = 0;
    int uIndex = frameSize;
    int vIndex = frameSize + chromaSize;

    int R, G, B, Y, U, V;
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            B = bgr[(i * width + j) * 3 + 0];
            G = bgr[(i * width + j) * 3 + 1];
            R = bgr[(i * width + j) * 3 + 2];

            //BGR to YUV
            Y = ((66 * R + 129 * G + 25 * B + 128) >> 8) + 16;
            U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128;
            V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128;

            yuv420p[yIndex++] = (unsigned char)((Y < 0) ? 0 : ((Y > 255) ? 255 : Y));
            if (i % 2 == 0 && j % 2 == 0)
            {
                yuv420p[uIndex++] = (unsigned char)((U < 0) ? 0 : ((U > 255) ? 255 : U));
                yuv420p[vIndex++] = (unsigned char)((V < 0) ? 0 : ((V > 255) ? 255 : V));
            }
        }
    }
}
//kernel function to convert bgr to yuv420sp
__global__ void bgr2yuv420sp(uchar3 *  d_in, unsigned char * d_out,
                               uint imgheight, uint imgwidth)
{

    int col_num = blockIdx.x*blockDim.x+threadIdx.x;
    int row_num = blockIdx.y*blockDim.y+threadIdx.y;

    if ((row_num < imgheight) && (col_num < imgwidth))
    {
//        uint32_t a = *((uint32_t *)&dinput[global_offset*3]);
        int global_offset = row_num*imgwidth+col_num;

        int r,g,b;
        r = int(d_in[global_offset].z);
        g = int (d_in[global_offset].y);
        b = int (d_in[global_offset].x);


        d_out[row_num * imgwidth + col_num] = bgr2y(r,g,b);
        if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
            int uv_offset = imgwidth*imgheight+((row_num>>1)*imgwidth)+col_num;
            d_out[uv_offset] = bgr2u(r,g,b);
            d_out[uv_offset+1] = bgr2v(r,g,b);

        }

    }
}
//kernel function to convert bgr to yuv420p
__global__ void bgr2yuv420p(uchar3 *  d_in, unsigned char * d_out,
                               uint imgheight, uint imgwidth)
{

    int col_num = blockIdx.x*blockDim.x+threadIdx.x;
    int row_num = blockIdx.y*blockDim.y+threadIdx.y;

    if ((row_num < imgheight) && (col_num < imgwidth))
    {
//        uint32_t a = *((uint32_t *)&dinput[global_offset*3]);
        int global_offset = row_num*imgwidth+col_num;

        int r,g,b;
        r = int(d_in[global_offset].z);
        g = int (d_in[global_offset].y);
        b = int (d_in[global_offset].x);


        d_out[row_num * imgwidth + col_num] = bgr2y(r,g,b);
        if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
            int u_offset = imgwidth*imgheight+((row_num>>1)*(imgwidth>>1))+(col_num>>1);
            d_out[u_offset] = bgr2u(r,g,b);
            int v_offset = u_offset+((imgheight>>1)*(imgwidth>>1));
            d_out[v_offset] = bgr2v(r,g,b);

        }
    }
}


int main(void)
{

    const uint imgheight = 1000;
    const uint imgwidth = 1500;

    //input and output
    uchar3 *d_in;
    unsigned char *d_out;
    uchar3 *idata = new uchar3[imgheight*imgwidth];
    unsigned char *odata = new unsigned char[imgheight*imgwidth*3/2];
    unsigned char *cdata = new unsigned char[imgheight*imgwidth*3/2];
    uchar3 pix;
    for (int i = 0; i < imgheight*imgwidth; i++){
      pix.x = (rand()%30)+40;
      pix.y = (rand()%30)+40;
      pix.z = (rand()%30)+40;
      idata[i] = pix;}
    for (int i = 0; i < imgheight*imgwidth; i++) idata[i] = pix;
    bgr_to_yuv420p(cdata, (unsigned char*) idata, imgwidth, imgheight);
    // malloc memo in gpu
    cudaMalloc((void**)&d_in, imgheight*imgwidth*sizeof(uchar3));
    cudaMalloc((void**)&d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2);

    //copy image from cpu to gpu
    cudaMemcpy(d_in, idata, imgheight*imgwidth*sizeof(uchar3), cudaMemcpyHostToDevice);

    dim3 threadsPerBlock(32, 32);
    dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x,
                       (imgheight + threadsPerBlock.y - 1) / threadsPerBlock.y);

    //run kernel function
    bgr2yuv420p<<<blocksPerGrid, threadsPerBlock>>>(d_in, d_out, imgheight, imgwidth);

    cudaDeviceSynchronize();

    //copy yuv420p from gpu to cpu
    cudaMemcpy(odata, d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2, cudaMemcpyDeviceToHost);
    for (int i = 0; i < (imgwidth*imgheight*3/2); i++) if (odata[i] != cdata[i]) {std::cout << "mismatch at: " << i << " was: " << (int)odata[i] << " should be: " << (int)cdata[i] << std::endl; return 0;}
    cudaFree(d_in);
    cudaFree(d_out);


    return 0;

}

मैं इस कोड या मेरे द्वारा पोस्ट किए गए किसी अन्य कोड के लिए शुद्धता का दावा नहीं करता। मेरे द्वारा पोस्ट किए गए किसी भी कोड का उपयोग करने वाला कोई भी व्यक्ति अपने जोखिम पर ऐसा करता है। मैं केवल यह दावा करता हूं कि मैंने मूल पोस्टिंग में पाई गई कमियों को दूर करने का प्रयास किया है, और उसका कुछ स्पष्टीकरण प्रदान करता हूं। मैं यह दावा नहीं कर रहा हूं कि मेरा कोड दोष-मुक्त है, या यह किसी विशेष उद्देश्य के लिए उपयुक्त है। अपने जोखिम पर इसका इस्तेमाल करें (या नहीं)।

रॉबर्ट क्रोवेल्ला अग्रिम धन्यवाद और कुछ देरी के लिए मुझे खेद है।

मेरा उद्देश्य bgr2yuv420p को कन्वर्ट करना है (bgr या rgb सिर्फ ऑर्डर है, इसे अनदेखा करें)। आपका कोड rgb2yuv420sp कन्वर्ट करना है।

यहाँ मेरा सीपीयू कोड है:

void bgr_to_yuv420p(unsigned  char* yuv420p, unsigned char* bgr, int width, int height)
{
    if (yuv420p == NULL || bgr== NULL)
        return;
    int frameSize = width*height;
    int chromaSize = frameSize / 4;

    int yIndex = 0;
    int uIndex = frameSize;
    int vIndex = frameSize + chromaSize;

    int R, G, B, Y, U, V;
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            B = bgr[(i * width + j) * 3 + 0];
            G = bgr[(i * width + j) * 3 + 1];
            R = bgr[(i * width + j) * 3 + 2];

            //BGR to YUV
            Y = ((66 * R + 129 * G + 25 * B + 128) >> 8) + 16;
            U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128;
            V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128;

            yuv420p[yIndex++] = (unsigned char)((Y < 0) ? 0 : ((Y > 255) ? 255 : Y));
            if (i % 2 == 0 && j % 2 == 0)
            {
                yuv420p[uIndex++] = (unsigned char)((U < 0) ? 0 : ((U > 255) ? 255 : U));
                yuv420p[vIndex++] = (unsigned char)((V < 0) ? 0 : ((V > 255) ? 255 : V));
            }
        }
    }
}

यहां आपका सीपीयू कोड है:

// I have no idea if these are the correct conversion formulas
// I simply lifted what I saw in your host code so that we 
// are using the same conversion calculations in host and device

__host__ __device__ unsigned char bgr2y(int R, int G, int B){
  int Y = ((66 * R + 129 * G + 25 * B + 128) >> 8) + 16;
  return (unsigned char)((Y<0)? 0 : ((Y > 255) ? 255 : Y));}
__host__ __device__ int bgr2u(int R, int G, int B){
  int U = ((-38 * R - 74 * G + 112 * B + 128) >> 8) + 128;
  return (unsigned char)((U<0)? 0 : ((U > 255) ? 255 : U));}
__host__ __device__ int bgr2v(int R, int G, int B){
  int V = ((112 * R - 94 * G - 18 * B + 128) >> 8) + 128;
  return (unsigned char)((V<0)? 0 : ((V > 255) ? 255 : V));}

void bgr_to_yuv420p(unsigned  char* yuv420p, unsigned char* bgr, int width, int height)
{
    if (yuv420p == NULL || bgr== NULL)
        return;
    int frameSize = width*height;

    int yIndex = 0;
    int uIndex = frameSize;

    int R, G, B;
    for (int i = 0; i < height; i++)
    {
        for (int j = 0; j < width; j++)
        {
            B = bgr[(i * width + j) * 3 + 0];
            G = bgr[(i * width + j) * 3 + 1];
            R = bgr[(i * width + j) * 3 + 2];

            //BGR to YUV
            yuv420p[yIndex++] = bgr2y(R,G,B);
            if (i % 2 == 0 && j % 2 == 0)
            {
                yuv420p[uIndex] = bgr2u(R,G,B);
                yuv420p[uIndex+1] = bgr2v(R,G,B);
                uIndex+=2;
            }
        }
    }
}

जाहिर है, यह अलग है। मेरे कोड में, u और v का अपना index.

    int frameSize = width*height;
    int chromaSize = frameSize / 4;
    int yIndex = 0;
    int uIndex = frameSize;
    int vIndex = frameSize + chromaSize;

आपके कोड में: यू और वी इंडेक्स साझा करते हैं।

int frameSize = width*height;
int yIndex = 0;
int uIndex = frameSize;

तो, यह अभी भी हल नहीं किया जा सकता है। पीएस: मैं ओपनसीवी एपीआई द्वारा परिवर्तित छवि के साथ अपनी परिणाम छवि दिखाने और तुलना करने के लिए ओपनसीवी का उपयोग करता हूं।