the launch timed out and was treminated. 大虾们快来分析解决！！！

[风停雨歇云淡]

代码如下：
__global__ static void GetkeyFramesNum(
float *frameangles,int *keyFramesNum,int lsize,float e,
int *pair ,int * sMark,int *oMark,int *lock,int *thread)
{
//初始化


extern __shared__ int number[];
const int tid=threadIdx.x;
float *allframesangles=frameangles;
int *trackMark=sMark;//一个线程的起始坐标 也是可以开始运行的标志 非零时
int *overMark=oMark;
int *threadnum=lock;
//初始化结束

//误差处理初始化

float error=e;

//误差处理初始化结束

//并行处理



if(tid==0)
{

int loop;
if(lsize>THREAD_NUMONE)loop=lsize;
else loop=THREAD_NUMONE;
for(int i=0;i<loop;i++)
{
if(i<THREAD_NUMONE)
{
trackMark[i]=0;
overMark[i]=0;
}
pair[i]=0;
}

trackMark[tid]=1;
pair[1]=lsize-1;
pair[lsize-1]=1;

overMark[tid]=lsize-1;
//threadnum[0]=1;

number[0]=2;//已取得关键帧数
number[1]=1;//其他线程结束空转的的标记
number[2]=1;//线程个数
number[3]=1;//互斥控制
threadnum[0]=1;//互斥控制
thread[0]=1;
thread[1]=1;


}
else
{
//是一些不可运行的代码做些无用功 相当于 sleep wait之类的
while(atomicExch(&trackMark[tid],1)==0)//如trackMark[tid]非零时停止
{
int i=1;
while(i<5)
{
i++;
}
if(trackMark[tid]==1)atomicExch(&trackMark[tid],0);

}
}
for ( ; trackMark[tid]!=overMark[tid]; )
{

int max_ind=0;//最大距离得帧id
float max_dist=0;//最大距离
float dist = 0;
for (int i = trackMark[tid]+1; i < pair[trackMark[tid]]; i++)
{
GetDistance(allframesangles,trackMark[tid],pair[trackMark[tid]],i,&dist);

if (dist > max_dist)
{
max_dist = dist;
max_ind = i;
}
}
float result=0;
GetVectorModel(allframesangles,trackMark[tid],pair[trackMark[tid]],&result);
float ratio=max_dist /result;
if (ratio > error)
{
while(atomicExch(&threadnum[0],0))
{
if(number[2]<THREAD_NUMONE)
{
overMark[tid]=max_ind;
atomicExch(&trackMark[1],max_ind);//下一线程的开始位置
overMark[1]=pair[trackMark[tid]];//下一线程的结束位置
number[2]++;

}
number[0]++;
thread[number[0]-1]=tid+1;
atomicExch(&threadnum[0],1);
break;
}
pair[max_ind] = pair[trackMark[tid]];
pair[trackMark[tid]] = max_ind;
}
else
{
trackMark[tid] = pair[trackMark[tid]];

}
}


__syncthreads();

if(tid==0)
{
int j=1;
keyFramesNum[0]=0;
for (int i = 1; i < lsize; i++)
{
//if (pair[i] != 0)
//{
keyFramesNum[j++]=pair[i];
//if(pair[i] != 0) printf("%d \n",i);

//}

}
}
}

程序逻辑：
1.首先由0号线程将一个数据分割为两段
2.再由0，1号线程分为四段
以此类推...
在这里只有0号分完1才执行。只有0，1再分完2，3才执行。。
我的程序写的在哪里出问题了？
运行结果：
Cuda error:Kernel execution failed
in file ‘.cu文件路径’in line 493: the launch timed out and was treminated.
这是在ctrl+F5下运行的结果。
直接debug 下提示代码段中存在异常：
CUDA2.exe 中的 0x75869617 处最可能的异常: Microsoft C++ 异常: 内存位置 0x002ff4fc 处的 cudaError_enum

而且在两种情况下显示器都会黑屏以下后提示显卡驱动停止工作 且恢复正常。

我的想法及疑问：
1.程序中分配的空间各线程都可以随意使用吗？
2.程序中的互斥操作对不对？
或者其他什么原因...
各位大虾帮忙解决。多多感谢！！！

[无心人_过过小日子]

gpu上不会主动切换线程的.所以,一般不建议用原子操作+while来实现同步等.
LZ的程序基本上是死锁了。

[风停雨歇云淡]

[Quote=引用 1 楼 l7331014 的回复:]
gpu上不会主动切换线程的.所以,一般不建议用原子操作+while来实现同步等.
LZ的程序基本上是死锁了。
[/Quote]

哦,那有没有其他方式实现互斥访问啊.对一段代码的及一个变量的.

[风停雨歇云淡]

[Quote=引用 1 楼 l7331014 的回复:]
gpu上不会主动切换线程的.所以,一般不建议用原子操作+while来实现同步等.
LZ的程序基本上是死锁了。
[/Quote]
对于代码段的互斥的我通过多个原子操作实现.是可行的吧?

一个问题是我如何让编号靠后的线程等待前面的线程为其分配开始标志后再继续运行下面的代码.

[风停雨歇云淡]

多谢提醒!此题已解决.换一种方式,大功告成.

[tengwl]

参照http://topic.csdn.net/u/20100125/09/542cbdc3-55be-4545-a52e-7c2f213c8a68.html