大侠帮我看一下, 我的内核程序

[yinquan8241]

// 这个函数确保data 的值在0~255之间
__device__ void ClipByte(short *data)
{
*data = (*data) < 0 ? 0 : ((*data) > 255 ? 255 : (*data));
}

__global__ static void GPU_YUV2RGB(BYTE *Y, BYTE *Cb, BYTE *Cr, BYTE *RGBbuf, int width, int height)
{
int i, j, k;
for(i=0, k=height-1; i<width, k>=0; i++, k--)
{
for(j=0; j<width; j++)
{
int w = (i / 2) * (width / 2) + j / 2;
short temp = -1;
int index_y = i * width + j;
int index_buf = k * width * 3 + j * 3;

// B通道
temp = Y[index_y] + 1.779 * (Cb[w] - 128);
ClipByte(&temp);
RGBbuf[index_buf] = temp;

// G通道
temp = Y[index_y] - 0.3456 * (Cb[w] - 128) - 0.7169 * (Cr[w] -128);
ClipByte(&temp);
RGBbuf(index_buf + 1) = temp;

// R通道
temp = Y[index_y] + 1.4075 * (cR[w] - 128);
ClipByte(&temp);
RGBbuf[index_buf + 2] = temp;
}
}
}
// 上面这个函数主要是想在GPU上实现YUV到RGB的转换.该算法的CPU版本运行良好.
现在问题是我一调用这个GPU版本的内核函数:
GPU_YUV2RGB<<<1, 1, 0>>>(gpu_Y, gpu_Cb, gpu_Cr, gpu_RGB, width, height);

cudaError_t err = cudaMemcpy(RGBbuf, gpu_RGB, sizeof(BYTE)*size, cudaMemcpyDevicetoHost);
拷贝出转换结果
程序运行到这一句的时候, err返回 cudaErrLaunchFailure.
前面的内存分配和拷贝都没问题 ,都返回了cudaSuccess
请大侠帮我看一下问题出在哪里??

[yinquan8241]

这里还没有做并行计算,只是做一个简单的内核调用

[无心人_过过小日子]

1)cudaErrLaunchFailure是内核失败.
因为内核是异步启动的.启动后,不等结束就执行到cudaMemcpy.(在执行cudaMemcpy前等待kernel结束).
因此,看上去好像是cudaMemcpy出错.
2)kernel中有bug(越界访问).LZ仔细看一下i循环中的结束条件对嘛?!

[无心人_过过小日子]

CPU版本运行良好....
结果正确?有疑问.

[yinquan8241]

[Quote=引用 3 楼 l7331014 的回复:]
CPU版本运行良好....
结果正确?有疑问.
[/Quote]
就是说
我的YUV转RGB函数
如果不写成内核函数的话，能运行成功。
我那个内核程序里的代码完全就是把CPU上的函数拷贝过去的

[无心人_过过小日子]

[Quote=引用 4 楼 yinquan8241 的回复:]
引用 3 楼 l7331014 的回复:
CPU版本运行良好....
结果正确?有疑问.

就是说
我的YUV转RGB函数
如果不写成内核函数的话，能运行成功。
我那个内核程序里的代码完全就是把CPU上的函数拷贝过去的
[/Quote]

说了:
LZ仔细看一下i循环中的结束条件对嘛?!

[yinquan8241]

楼上的大侠
i循环的结束条件是对的

[无心人_过过小日子]

[Quote=引用 6 楼 yinquan8241 的回复:]
楼上的大侠
i循环的结束条件是对的
[/Quote]

i<widt还是 i<height?

[OpenHero]

可以断定是内存越界了~
这其实sdk里面有个demo是nv12toagb的，这个例子可以让你来修改yuv2rgb，以前写过一个，不过哪个是商用的，不能公开给lz~~
lz可以仿照NV12ToARGB_drvapi.cu 来修改

[yinquan8241]

节后我再结帖
多谢大侠指点

[yinquan8241]

[Quote=引用 8 楼 openhero 的回复:]
可以断定是内存越界了~
这其实sdk里面有个demo是nv12toagb的，这个例子可以让你来修改yuv2rgb，以前写过一个，不过哪个是商用的，不能公开给lz~~
lz可以仿照NV12ToARGB_drvapi.cu 来修改
[/Quote]
没想到业界鼎鼎大名的开勇大哥也给我回帖了
感动
非常感动
节后结帖给分