मुझे छवि को bgr से yuv420p में बदलने की आवश्यकता है और ऐसा करने के लिए मैं सबसे पहले OpenCV का उपयोग करता हूं।

Mat img = imread("1.bmp");
Mat yuvImg;
cvtColor(img,yuvImg,COLOR_BGR2YUV_I420);

इसका परिणाम सामान्य है। हालांकि मेरी छवि बहुत बड़ी है और इसका पिक्सेल लगभग 6400 * 2000 है। मुझे लगता है कि opencv api cvtcolor के साथ bgr को yuv420p में बदलने में बहुत अधिक समय खर्च होता है।

फिर मैं इसे स्वयं परिवर्तित करने का निर्णय लेता हूं और इसे कूडा से गति देता हूं।

यह cpu में कोड है:

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));
            }
        }
    }
}

मैं bgr_to_yuv420p(...) कोड का परीक्षण करता हूं और परिणाम भी सामान्य है।

फिर मैं इसे क्यूडा के साथ तेज करता हूं।

यहां मेरे सभी कोड में कर्नेल फ़ंक्शन और टेस्ट फ़ंक्शन शामिल हैं।

#include <iostream>
#include <time.h>
#include <vector_types.h>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include "opencv2/highgui.hpp" 
#include "opencv2/opencv.hpp"
using namespace cv;
using namespace std;

//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] = ((66*r + 129*g + 25*b) >> 8) + 16;
        if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
            int uv_offset = imgwidth*imgheight+((row_num*imgwidth))+col_num;
            d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
            d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;

        }

    }
}

int main(void)
{

    Mat srcImage = imread("1.bmp");
    imshow("srcImage", srcImage);
    const uint imgheight = srcImage.rows;
    const uint imgwidth = srcImage.cols;

    Mat nv12Image(imgheight * 3 / 2, imgwidth, CV_8UC1, Scalar(255));

    //input and output 
    uchar3 *d_in;
    unsigned char *d_out;

    // 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, srcImage.data, 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(nv12Image.data, d_out, imgheight*imgwidth*sizeof(unsigned char) * 3 / 2, cudaMemcpyDeviceToHost);

    imshow("nv12",nv12Image);
    imwrite("cuda.bmp",nv12Image);

    cudaFree(d_in);
    cudaFree(d_out);


    return 0;

}

क्यूडा के साथ कोड चल सकता है लेकिन परिणाम सामान्य नहीं है। YUV420p का Y सामान्य है लेकिन U और V में कुछ गड़बड़ है। मुझे लगता है कि इसका कारण __global__ void bgr2yuv420p(...) में है।

if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
                int uv_offset = imgwidth*imgheight+((row_num*imgwidth))+col_num;
                d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
                d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;

            }

मैं बहुत कोशिश करता हूं लेकिन फिर भी इसे हल नहीं कर पाता। और मुझे rgb को yuv420p में बदलने के बारे में बहुत कम कोड मिलते हैं, अधिक कोड yuv420p को rgb में बदलने के बारे में हैं। तो मैं जानना चाहता हूं कि क्या कोई इसी तरह के सवाल का सामना कर रहा है या मुझे कुछ सलाह दे रहा है?

धन्यवाद रॉबर्ट क्रोवेल्ला। यह रहा मेरा अपडेट-1

मैं रॉबर्ट क्रोवेला की सलाह का पालन करता हूं और कर्नेल फ़ंक्शन को इस तरह बदलता हूं:

//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] = ((66*r + 129*g + 25*b) >> 8) + 16;
            if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
                int uv_offset = imgwidth*imgheight+((row_num>>1)*imgwidth)+col_num;
                d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
                d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;

            }

        }
    }

मैं उत्साह के साथ नए कर्नेल का परीक्षण करता हूं, लेकिन परिणाम भी सामान्य नहीं है। अद्यतन कर्नेल फ़ंक्शन के साथ मेरी परिणाम छवि यहां दी गई है। yuv420p इमेज मैंने खुद बदली है

फिर ओपनसीवी एपीआई द्वारा परिवर्तित सामान्य परिणाम छवि यहां है। yuv420p इमेज को opencv api द्वारा रूपांतरित किया गया

जैसा कि हम देख सकते हैं, दो छवियों के बीच का अंतर यू और वी है। मैंने कर्नेल फ़ंक्शन में यू और वी की अनुक्रमणिका पहले ही बदल दी है, यानी।

if(((threadIdx.x & 1) == 0)  && ((threadIdx.y & 1) == 0)){
                int uv_offset = imgwidth*imgheight+((row_num >>1)*imgwidth)+col_num;
                d_out[uv_offset] = ((112*r + -94*g + -18*b) >> 8) + 128;
                d_out[uv_offset+1] = ((-38*r + -74*g + 112*b) >> 8) + 128;

            }

मुझे लगता है कि यह काम करेगा लेकिन ऐसा नहीं है। कोई अन्य सलाह? रॉबर्ट क्रोवेल्ला

संपादित करें: समाधान रॉबर्ट क्रोवेल्ला का नवीनतम उत्तर है। मैंने इसे दोबारा जांच लिया है और यह वास्तव में सही है।

-1
yuyuyu 27 अप्रैल 2020, 13:54
आपकी कर्नेल परिभाषा bgr2yuv420p है, लेकिन आपका कर्नेल लॉन्च rgb2yuv420p है, इसलिए आपका कोड संकलित नहीं होगा। मुझे लगता है कि यह वास्तव में वह कोड नहीं है जिसे आप चला रहे हैं। आपके कर्नेल में, uv_offset की आपकी गणना बाद की पंक्तियों में अवैध, सीमा से बाहर पहुंच उत्पन्न कर रही है। आप शायद उस अनुक्रमण का अध्ययन करना चाहें और यहां वर्णित विधि का उपयोग करना चाहें। डिबग में मदद करें।
 – 
Robert Crovella
27 अप्रैल 2020, 17:54
अपनी uv_offset गणना के बारे में ध्यान से सोचें। imgwidth*imgheight ऑफसेट आपको Y क्षेत्र (सही ढंग से) से आगे ले जाता है, लेकिन उस बिंदु से, क्या यूवी प्लानर क्षेत्र में पंक्ति द्वारा अनुक्रमित करने के लिए row_num*imgwidth का उपयोग करना सही है? (संकेत, ऐसा नहीं है। आपके पास यूवी प्लानर क्षेत्र में इतनी पंक्तियाँ नहीं हैं, आपके पास केवल आधी पंक्तियाँ हैं)।
 – 
Robert Crovella
27 अप्रैल 2020, 18:00
मुझे लगता है कि अगर आप यह बदलाव करते हैं तो आप कुछ सही के करीब पहुंच जाएंगे: int uv_offset = imgwidth*imgheight+(((row_num>>1)*imgwidth))+col_num;
 – 
Robert Crovella
27 अप्रैल 2020, 18:09
1
आपके bgr_to_yuv420p CPU कोड की तुलना आपके bgr2yuv420p GPU कर्नेल से करने पर हम यह भी देखते हैं कि U और V स्टोरेज का क्रम उलट गया है, और कुछ अन्य गणना अंतर हैं। अगर मैं उन सभी को छाँटता हूँ, तो मुझे उनके बीच मेल खाने वाले परिणाम मिल सकते हैं। ध्यान दें कि आपका bgr_to_yuv420p CPU कोड Y, U, V के लिए प्लानर स्टोरेज को दर्शाता है जबकि आपका GPU कोड Y प्लेन और UV इंटरलीव्ड प्लेन के लिए सेमी-प्लानर स्टोरेज डिलीवर कर रहा है।
 – 
Robert Crovella
27 अप्रैल 2020, 18:43
कर्नेल फ़ंक्शन rgb2yuv420p() और bgr2yuv420p() के बारे में निम्न स्तर की त्रुटि के लिए क्षमा करें। कर्नेल फ़ंक्शन चलाते समय मैं बस गलत इनपुट करता हूं। वास्तव में, मैं bgr2yuv420p() को परिभाषित करता हूं और मैं bgr2yuv420p() भी चलाता हूं। मैं इसके बारे में दोबारा जांच करूंगा और अपना प्रश्न संपादित करूंगा।
 – 
yuyuyu
28 अप्रैल 2020, 04:49

2 जवाब

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

  • आपके सीपीयू और जीपीयू कोड के बीच आर, जी, बी को वाई, यू, वी में बदलने की गणना समान नहीं है। हाँ, यह मायने रखता है।
  • आपके सीपीयू कोड में प्लानर वाई, यू, वी स्टोरेज है। इसका मतलब है कि वाई का अपना विमान है, यू का अपना विमान है, और वी का अपना विमान है। आपका 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;

}

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

1
Robert Crovella 12 मई 2020, 20:24
मैं आपके कोड का परीक्षण करता हूं और परिणाम अच्छा है: सीपीयू में परिणाम और जीपीयू में परिणाम पूरी तरह से समान है। लेकिन आपका कोड मेरे से अलग है। मैं yuv420p का उपयोग करता हूं और आप yuv420sp का उपयोग करते हैं। मैं नीचे इसके बारे में चर्चा करते हुए एक उत्तर प्रदान करूंगा।
 – 
yuyuyu
12 मई 2020, 09:36
आपके प्रश्न में, आपके पास एक सीपीयू कोड है जो प्लानर है और एक जीपीयू कार्यान्वयन जो सेमी-प्लानर है। बेशक परिणाम मेल नहीं खा सकते। परिणाम संगठन अलग है। मुझे समकक्ष आउटपुट दिखाने के लिए एक दृष्टिकोण या दूसरे को चुनना पड़ा, मैंने सेमी-प्लानर चुना। एक GPU कोड होना पूरी तरह से संभव है जो एक प्लानर प्रारूप परिणाम को आउटपुट करता है।
 – 
Robert Crovella
12 मई 2020, 18:25
मैंने प्लानर स्टोरेज जेनरेट करने के लिए अपना कोड अपडेट कर लिया है।
 – 
Robert Crovella
12 मई 2020, 20:24
बहुत बढ़िया! आपका नया कोड काम करता है। मैं वास्तव में आपकी सराहना करता हूं। आप मेरे प्रश्न का उत्तर समय पर देते रहें और अंत में उसे हल करते रहें। मैं फिर से धन्यवाद व्यक्त करना चाहता हूं। आशा है कि यह उत्तर किसी और के लिए सहायक होगा।
 – 
yuyuyu
13 मई 2020, 12:33

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

मेरा उद्देश्य 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;

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

0
yuyuyu 12 मई 2020, 10:44