unknown error

加速计算 CUDA
1

运行到一半，程序突然退出，报错cuda error in file"C:/…"in line 252:unknown error.

代码部分如下：
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result1, width, height, dLeftGradient);
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result2, width, height, dRightGradient);
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result3, width, height, dUpGradient);
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result4, width, height, dDownGradient);

float *hresult1;

hresult1 = new float[blockCount];
float *hresult2;
hresult2 = new float[blockCount];
float *hresult3;
hresult3 = new float[blockCount];
float *hresult4;
hresult4 = new float[blockCount];
for(int i =0;i<blockCount;i++)
{
hresult1[i] = 0;
hresult2[i] = 0;
hresult3[i] = 0;
hresult4[i] = 0;
}
CUDA_SAFE_CALL( cudaThreadSynchronize() );

运行到CUDA_SAFE_CALL( cudaThreadSynchronize() )就突然退出报错。有高手知道原因吗？

2

这个错误一般在两种情况下会出现。
(1)你人为的让你的device code挂掉，比如你嵌入了不正确的ptx汇编，或者干脆你故意，如asm(“trap;”);
(2)你的device code里面使用了非法指针。比如你的__global__里面尝试访问无效位置int *p=0; *p=1234;就会立刻出现unknown error而挂掉。

最快的解决方案：立刻检查你的device code, 仔细观察指针(包括数组名)，找出使用不当的地方。

另：你这是在做什么？
for(int i =0;i<blockCount;i++)
{
hresult1 = 0;
hresult2 = 0;
hresult3 = 0;
hresult4 = 0;
}

3

谢谢啊~~代码calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result1, width, height, dLeftGradient)是计算数组dLeftGradient中的最大值返回给result1，result1也是一个数组，元素个数为block数（类似树状加法优化），每个元素为一个block中数字的最大值，然后这样计算出来四个result数组，再求四个result数组的最大值，所有需要再进行4个result数组数字的比较。
我想把result1到result4传到CPU做这个工作，但是执行到上面的地方就出错。那个循环语句我贴错了，
for(int i =0;i<blockCount;i++)
{
hresult1[i] = 0;
hresult2[i] = 0;
hresult3[i] = 0;
hresult4[i] = 0;
}
应该是这样，在CPU上实现数组初始化。我的device code是这样的：
printf(“hello cuda!”);
CUDA_SAFE_CALL( cudaThreadSynchronize() );
CUDA_SAFE_CALL(cudaMemcpy(hresult1,result1,sizeof(float)* blockCount, cudaMemcpyDeviceToHost));
CUDA_SAFE_CALL(cudaMemcpy(hresult2,result2,sizeof(float)* blockCount, cudaMemcpyDeviceToHost));
CUDA_SAFE_CALL(cudaMemcpy(hresult3,result3,sizeof(float)* blockCount, cudaMemcpyDeviceToHost));
CUDA_SAFE_CALL(cudaMemcpy(hresult4,result4,sizeof(float)* blockCount, cudaMemcpyDeviceToHost));

都是自带函数啊，怎么会出错呢。。（打印hello cuda后就退出）

4

汗，，贴过来不支持。。。。hresult1【i】这样的

5

汗，过一会跑又变结果了，原代码没动，还是到那个位置，同步和mencpy的地方然后显示器黑屏两秒然后报错the launch timed out and was treminated.还有显卡驱动程序停止工作并已经恢复。。。。

6

没有看懂你在说什么。2楼的回复的帖子。你照着检查了吗？
在一个context里面, 如果你一个调用出错，那么你会连续出错，无论下面是什么。也就是说，比如：你的一个launch失败了，那么它之后的所有的调用（比如memcpy, sync, 等等），都会返回这个错误的。
现在很明显你的设备代码里面有问题，如同2楼所说，你必须去检查一下。

7

你好~我把之前device里面的的 extern shared int 改为 shared float shared[1024]，然后没有报错unknown error了，但是在CUDA_SAFE_CALL(cudaMemcpy(hresult1,result1,sizeof(float)* blockCount, cudaMemcpyDeviceToHost))的时候报错the launch timed out and was treminated，然后注释掉这些Memcpy就能继续往下运行。。。不知道这又怎么回事啊。。。

8

[
每个调用都可能返回上个调用的错误的。你自己检查你的kernel是不是过长时间执行导致。这个错误和返回它的memcpy无关。

9

谢谢哥们。我也看有说kernel运行时间过长。。。但是这之前我就八个kernel：
calcLeftGradientD<<<blockCount,BLOCK_SIZE>>>(dLeftData, width, height, dLeftGradient);
calcRightGradientD<<<blockCount,BLOCK_SIZE>>>(dLeftData, width, height, dRightGradient);
calcUpGradientD<<<blockCount,BLOCK_SIZE>>>(dLeftData, width, height, dUpGradient);
calcDownGradientD<<<blockCount,BLOCK_SIZE>>>(dLeftData, width, height, dDownGradient);

calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result1, width, height, dLeftGradient);
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result2, width, height, dRightGradient);
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result3, width, height, dUpGradient);
calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result4, width, height, dDownGradient);

其中calcMaxGradientD<<<blockCount,BLOCK_SIZE>>>(result4, width, height, dDownGradient)是我根据深入浅出里面的例程改的比较计算数组的最大元素，具体代码如下：
global static void calcMaxGradientD(float *result,int width,int height,unsigned char *dGradient)
{
shared int shared[1024];
int i;

shared[threadIdx.x] = 0;
for(i = blockIdx.x * blockDim.x + threadIdx.x;i < width*height;i += blockIdx.x * blockDim.x)
{
shared[threadIdx.x] = dGradient[blockIdx.x * blockDim.x + threadIdx.x];
}
__syncthreads();

if(threadIdx.x < 512) 
{
	if(shared[threadIdx.x] < shared[threadIdx.x + 512])
		shared[threadIdx.x] = shared[threadIdx.x + 512];
}

__syncthreads();
if(threadIdx.x < 256)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 256])
shared[threadIdx.x] = shared[threadIdx.x + 256];
}
__syncthreads();
if(threadIdx.x < 128)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 128])
shared[threadIdx.x] = shared[threadIdx.x+128];
}
__syncthreads();
if(threadIdx.x < 64)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 64])
shared[threadIdx.x] = shared[threadIdx.x+64];
}
__syncthreads();
if(threadIdx.x < 32)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 32])
shared[threadIdx.x] = shared[threadIdx.x+32];
}
__syncthreads();
if(threadIdx.x < 16)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 16])
shared[threadIdx.x] = shared[threadIdx.x+16];
}
__syncthreads();
if(threadIdx.x < 8)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 8])
shared[threadIdx.x] = shared[threadIdx.x+8];
}
__syncthreads();
if(threadIdx.x < 4)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 4])
shared[threadIdx.x] = shared[threadIdx.x + 4];
}
__syncthreads();
if(threadIdx.x < 2)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 2])
shared[threadIdx.x] = shared[threadIdx.x+2];
}
__syncthreads();
if(threadIdx.x < 1)
{
if(shared[threadIdx.x] < shared[threadIdx.x + 1])
shared[threadIdx.x] = shared[threadIdx.x + 1];
}
__syncthreads();

if(threadIdx.x == 0)
{
result[blockIdx.x] = shared[0];
}
}
这个kernel我按照例程改写的，运算时间也应该不多啊，数组有480*272个数。请指教。

10

看写字板太费力了。。sigh。。
貌似数组不是很大，用gpu不一定有cpu的release优化过的快。。

[ 本帖最后由 wscuiqiu 于 2010-11-2 17:42 编辑 ]