我的问题是：我有压缩好MPEG视频源，已用cudaSDK中自带的cudaDecodeGL解码解码成NV12格式的数据，而我需要的是RGB24格式的数据。
SDK中的demo：cudaDecodeGL工程中自带了cuda文件NV12ToARGB_drvapi.cu，可以完成对视频解码后数据的后处理，处理的结果是NV12转ARGB32。
我想要的是RGB24，现在又两个方法：
一：接着现有的程序，先将NV12转成ARGB，然后再将ARGB转成RGB，最好在GPU中进行；
二：修改现有程序NV12ToARGB_drvapi.cu，将NV12直接转成RGB，也是在GPU中进行。
注： 由于解码程序是用的cuda驱动API写的，又驱动API和运行API不能同时应用于一个工程中，所以整个数据转化部分也必须用驱动API写。
各位，请问谁那有改好的程序或者给我指点一下也行的 万分谢谢！！！
楼主很抱歉，我不懂OpenGL, 建议其他版主/会员/官方支持为您服务。
对于您的部分问题：“注： 由于解码程序是用的cuda驱动API写的，又驱动API和运行API不能同时应用于一个工程中，所以整个数据转化部分也必须用驱动API写”
我想说下我的观点：driver api可以和runtime api混用的，楼主。
（1）简单的您可以先尝试任意runtime api的能lazy initialize出context的函数(例如上去先来一个cudaMalloc)，然后照旧使用您的driver api(包括module载入，kernel查找，设置参数和调用）。
（2）如果您在runtime api之前已经初始化过context, 或者其他原因不能/不愿和runtime api的隐式的context混用的话，您可以考虑下反复的cuCtxPush/PopCurrent()来保存和切换您的多个context.
我想这对您应该很简单。不妨尝试下？
我想你可以用runtime api写一个kernel来对ARPG转成RPG的格式是没有问题的！
2楼版主给出的driver api和runtime api混用是可行的，前提是你要在工程中包含正确的头文件以及链接正确的库
首先谢谢您的解答，关于驱动API和运行API是否可以混用的问题，我虽然没有照您的方法试验过，但是之前试过将运行API写的核函数用于工程中，但总是提示错误：
error LNK2019: 无法解析的外部符号 _run_yuv2rgb_420，该符号在函数 “public: bool __thiscall Ccamera_guiDlg::YUV420_to_RGB(void)” (?YUV420_to_RGB@Ccamera_guiDlg@@QAE_NXZ) 中被引用
run_yuv2rgb_420是.cu文件中的一个函数，也已经用了exturn标明是外部函数，但就是不识别，请问楼主有遇见过类似的情况吗？
另外，您的建议我看了，你只是要我实验一个函数cudaMalloc()，很明显这个函数是运行API写的，但是会不会因为是cuda自带的函数库，里面已经做了处理，而能和驱动API混用，只是咱们不知道而已。
不过，还是谢谢楼主，我会试着按照您的建议试验一下的
我试过了用runtime API写一个kernel，但是总是报错 error LNK2019: 无法解析的外部符号 _run_yuv2rgb_420，该符号在函数 “public: bool __thiscall Ccamera_guiDlg::YUV420_to_RGB(void)” (?YUV420_to_RGB@Ccamera_guiDlg@@QAE_NXZ) 中被引用。
要是按照您说的，可以通过加一些头文件及连接库解决我的问题，就太好了。
我是迫于项目需要，才接触cuda的，实属菜鸟一枚，还望高人给我指点一下的
programming guide 3.4指出了，runtime api和driver api可以混用，并且如果是以driver api建立的上下文，随后调用的runtime api将直接在driver api建立的上下文运行程序。
你这个报错是因为你调用的函数编译器找不到…
我想你可以尝试以下方式解决：
将你写的kernel单独放在一个.cu文件中
文件形式是：
#include “cuda.h”
#ifndef XXXX
#define XXXX
global void static kernel(…)
{…}
#endif
然后再你要调用这个函数的文件直接包含这个.cu文件
可能跟您说的差不多，我的程序也是放在了yuv420_rgb.cu中，
#pragma once
include <cutil_inline.h>
include <cutil_math.h>
include"video_info.h"
global void yuv2rgb_420_8lu(unsigned char * in_data,unsigned char *out_data)
{。。。。。。。。}
extern “C”
void run_yuv2rgb_420(unsigned char * in_data,unsigned char *out_data)
{
dim3 dimGrid(DIV_UP(VWIDTH, 16), DIV_UP(VHEIGHT, 16));
dim3 dimBlock(16,16);
yuv2rgb_420_8lu<<<dimGrid,dimBlock>>>(in_data,out_data);
}
跟您不一样的地方应该是
文件形式是：
include “cuda.h”
#ifndef XXXX
define XXXX
楼主您好，看到您的回帖了。
您的报错可能有如下原因：
（1）您的kernel没参与编译。
（2）您的kernel参与了编译，但没参与链接。
建议的解决方案：
检查是否参与编译和连接，（或者干脆使用楼上tianyuan2008方式）。
感谢继续回复，根据您最新发表的文章，此原因可能是因为C++对函数进行的name mangling导致的，
您的run_yuv2rgb_420, 是作为C函数编译的，即没有使用c++ name mangling。如果您的调用者的看到的声明里没有extern "C"则可能无法连接。
所以，您的声明，然后被别的函数、文件、obj使用的时候，别忘记声明上extern "C"了。建议按此修正。
祝您调试愉快！
好的 回头我好好检查的
以您的经验，你觉得如果是我的kernel没参与编译或者参与了编译，但没参与链接，最有可能是什么造成的？
不好意思。有点混乱。根据链接器的记录里的命名格式(_然后加C函数名), 您的确是extern "C"这样声明了。请排除此原因。
那可能是没有参与编译和连接。为何没有参与我就不知道了，如果您用VS, 可以看下您的.cu是否是CUDA C/C++类型？如果是的话。nv提供的build customazations/rules会进行编译和链接的。
祝您一切顺利！
我的.cu大概是下面这个样子
#pragma once
include <cutil_inline.h>
include <cutil_math.h>
include"video_info.h"
//========YUV420转RGB 以R G B顺序存放/////////////////////////////////////
global void yuv2rgb_420_8lu(unsigned char * in_data,unsigned char *out_data)
{
//获取像素索引 //获得线程ID
int x = blockIdx.x * (blockDim.x << 1) + (threadIdx.x << 1);
int y = blockIdx.y * (blockDim.y<<1) + (threadIdx.y<<1);
if (x >= VWIDTH)
return; //i = 720 - 1;
if (y >= VHEIGHT)
return; // j = 384 - 1;
//v1
float Y = in_data[y*VWIDTH + x]; //4个Y共用一个U和V
float U = in_data[Y_PLANE + UV_WIDTH*(y>>1) + (x>>1)];
float V = in_data[Y_PLANE + UV_PLANE + UV_WIDTH*(y>>1) + (x>>1)];
unsigned char* pOut = &out_data[3*VWIDTH*y+6*(x>>1)];
*pOut++ = (unsigned char)clamp(Y + (1.4075f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y - (0.3455f * (U - 128.f) + 0.7169f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y + (1.7790f * (U - 128.f)), 0.f, 255.f);
Y = in_data[y*VWIDTH + x+1];
*pOut++ = (unsigned char)clamp(Y + (1.4075f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y - (0.3455f * (U - 128.f) + 0.7169f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y + (1.7790f * (U - 128.f)), 0.f, 255.f);
Y = in_data[(y+1)*VWIDTH + x];
pOut = &out_data[3*VWIDTH*(y+1)+6*(x>>1)];
*pOut++ = (unsigned char)clamp(Y + (1.4075f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y - (0.3455f * (U - 128.f) + 0.7169f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y + (1.7790f * (U - 128.f)), 0.f, 255.f);
Y = in_data[(y+1)*VWIDTH + x+1];
*pOut++ = (unsigned char)clamp(Y + (1.4075f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y - (0.3455f * (U - 128.f) + 0.7169f * (V - 128.f)), 0.f, 255.f);
*pOut++ = (unsigned char)clamp(Y + (1.7790f * (U - 128.f)), 0.f, 255.f);
//v2 /*一下内容跟v1类似
//v3
//v4
//v5
//v6
//v7
//v8
}
extern “C”
void run_yuv2rgb_420(unsigned char * in_data,unsigned char *out_data)
{
dim3 dimGrid(DIV_UP(VWIDTH, 16), DIV_UP(VHEIGHT, 16));
dim3 dimBlock(16,16);
//execute the kernel
yuv2rgb_420_8lu<<<dimGrid,dimBlock>>>(in_data,out_data);
}
感谢您的补充。不过，请参考下11楼的处理建议。
真是好东西呀
楼上，这种回复希望不要出现在此版块哦！
超版说的对！建议楼主下次不要这样发帖。