1、cuda计时器cudaEvent_t:代码如下` //cuda计算运行时间
float time_t=0;//cuda运行时间
cudaEvent_t start_t,stop_t;//开始和结束时间
//创建
cudaEventCreate(&start_t);
cudaEventCreate(&stop_t);
//记录cuda时间
cudaEventRecord(start_t,0);//记录
//运行函数
cudaEventRecord(stop_t,0);//记录
cudaEventSynchronize(start_t);
cudaEventSynchronize(stop_t);
cudaEventElapsedTime(&time_t,start_t,stop_t);//计算时间,单位ms
cudaEventDestroy(start_t);
cudaEventDestroy(stop_t);
2、_syncthreads
_syncthreads()是cuda的内建函数,用于块内线程通信。
保证block内的所有线程都已经运行到调用_syncthreads()的位置。
(1)、块内线程同步
__share__ val[];
...
if(index < n)
{
if(tid condition)
{
do something with val;
}
__syncthreads();
do something with val;
__syncthreads();
}
(2)、
`share val[];
…
if(index < n)
{
if(tid condition)
{
do something with val;
__syncthreads();
}
else
{
do something with val;
__syncthreads();
}
}
这种结构将块内线程分成两部分,每一部分对共享存储器进行些操作,并在各自部分里同步.这种结构空易出现的问题是若两部分都要对某一地址的共享存储器进行写操作,将可能出
现最后写的结果不一致错误.要让错误不发生需要使用原子操作.
(3)、
__share__ val[];
....
if(index < n)
{
if(tid condition)
{
do something with val;
__syncthreads();
}
do something with val;
}
这种结构,块内只有部分线程对共享存储器做处理,并且部分线程是同步.那些不满足if条件的线程,会直接执行后面的语句.若后面的语句里面和if里面的语句都对共享存储器的同一
地址进行写操作时将会产生wait forever。若没有这种情况出现,程序则可以正常执行完.
在使用if condition 和__syncthreads(),最好使用第一结构,容易理解,不容易出错~