time.h 与cudaEvent记录的时间上差距很大不知问题所在

system · 2013 年5 月 16 日 13:13

本菜鸟刚刚入门cuda不久，再评测自己程序性能的时候，发现使用cudaevent 与调用 time.h去记录时间，发现两者差距很大，不知道是我自己使用上有问题，还是认识理解上存在差距。付上使用代码：

clock_t ss,eend;
ss=clock(); 
eend=clock();
printf("Total Time=%lf\n",(double)(eend - ss)/CLOCKS_PER_SEC);

float elapsedTime; 
cudaEvent_t start, stop;

cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);

cudaEventRecord(stop,0);
cudaEventSynchronize(stop); 
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("%f\n",elapsedTime);
cudaEventDestroy(start);
cudaEventDestroy(stop);

中间插入的是自己的kenel函数，对于同一个kenel函数结果如下：

Time.h Time=0.001000
cudaevent time=0.001376

Time.h Time=0.001000
cudaevent time=0.715872

system · 2013 年5 月 16 日 13:35

LZ您好，这个建议使用event来计时的。

以及在host端使用clock()函数不很推荐，具体的一些情况您可以参考一下下帖：
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6613&extra=&page=1

以及您给出的计时数据说明了同一个kernel两次运算时间不同么？原文看的好像不是很清楚，请LZ再详细说明一下。

大致如此，祝您好运~

system · 2013 年5 月 16 日 13:44

谢谢回答…第一个时间是指一个问题在cpu上完全运行的时间是多少，第二个时间是指在分配成kenel函数以后，在gpu上运行的时间。

根据您的回答，一般而言cuda中用cudaevent更加准确是吗？还有为什么我的程序会在gpu中运行中比cpu中慢很多？慢了差不多10倍…感觉kenel的优化也进行过一些了

system · 2013 年5 月 16 日 13:54

LZ您好，推荐使用event计时的。

以及您指出：“第二个时间是指在分配成kenel函数以后，在gpu上运行的时间。 ”您给的两组数据是同一个kernel函数使用同样的参数么？

以及，“还有为什么我的程序会在gpu中运行中比cpu中慢很多？慢了差不多10倍…感觉kenel的优化也进行过一些了”，这个和您的算法实现，计算规模，数据传输等等因素有关的。有些算法不适合GPU并行处理的，有些计算量太小，CPU完成计算的时间甚至比给GPU传递数据的时间更短。
所以，这个和具体情况有关，无法直接指出是何原因。但有一点可以确定，就是GPU程序不一定能比CPU快。

大致如此，供您参考。

祝您编码顺利~

system · 2013 年5 月 16 日 14:05

kernel函数中和cpu中运行的程序完全一样，就是对比同一个问题在两个地方不同的地方的运行情况。

实际我要处理的问题就是给10万个点，求取其中最小与原点的距离。那么实际就是最典型的归约问题，我完全按照书本的指示做的，附上代码，求版主指点。

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

const int N =1024*100; // 数组元素的个数
const int threadsPerBlock = 256;
const int blocksPerGrid = (N+threadsPerBlock-1) / ( 4*threadsPerBlock ) ;

__global__ void kenel( int2 *a, int *b)	 
{
	extern __shared__ int sdata[];
	int tx  = threadIdx.x;	
	int tid = blockIdx.x * blockDim.x + tx;
	int t=blockDim.x*blocksPerGrid;
	
	//每一个线程计算4个点，间隔就是所有线程的数量
	int minn=1000*1000;
	int2 a0;

	a0=a[tid]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;
	
	a0=a[tid+t]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;

	a0=a[tid+2*t]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;

	a0=a[tid+3*t]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;	
	
	sdata[tx]=minn;
	__syncthreads();

	
	if (tx < 128 &&  sdata[tx]> sdata[tx + 128]) sdata[tx]=sdata[tx+128];
	__syncthreads();
	if (tx <  64 &&  sdata[tx]> sdata[tx+ 64]) sdata[tx]=sdata[tx+64];
	__syncthreads();

	if (tx < 32)
	{
			if(sdata[tx]> sdata[tx + 32]) sdata[tx] = sdata[tx + 32]; 
			if(sdata[tx]> sdata[tx + 16]) sdata[tx] = sdata[tx + 16]; 
			if(sdata[tx]> sdata[tx + 8])  sdata[tx] = sdata[tx + 8]; 
			if(sdata[tx]> sdata[tx + 4])  sdata[tx] = sdata[tx + 4]; 
			if(sdata[tx]> sdata[tx + 2])  sdata[tx] = sdata[tx + 2]; 
			if(sdata[tx]> sdata[tx + 1])  sdata[tx] = sdata[tx + 1]; 			
	}
	
	if(tx==0) b[blockIdx.x]=sdata[0];
	return;
}

int main()
{	
	int2 *a;
	int *b;
	int2 *dev_a;
	int *dev_b;	
	int sharedMemSize=threadsPerBlock*sizeof(int);
	FILE *fp;

	//Cpu上空间分配
	a = (int2*)malloc( N*sizeof(int2) );
	b = (int*)malloc( blocksPerGrid*sizeof(int) );

	//Gpu上空间分配
	cudaMalloc( (void**)&dev_a, N*sizeof(int2) ) ;
	cudaMalloc( (void**)&dev_b, blocksPerGrid*sizeof(int) ) ;

	for(int i=0;i<N;i++)
		a[i].x=rand()%1000,a[i].y=rand()%1000;
	  
	//拷贝输入数据到GPU
	cudaMemcpy( dev_a, a, N*sizeof(int2), cudaMemcpyHostToDevice ) ;
		
	kenel<<<blocksPerGrid,threadsPerBlock,sharedMemSize>>>( dev_a, dev_b );

	//数据由GPU拷贝到CPU
	cudaMemcpy( b, dev_b,blocksPerGrid*sizeof(int), cudaMemcpyDeviceToHost ) ;
		
	int mindis=1000*1000;
	for(int i=0;i<blocksPerGrid;i++)
		if(b[i]<mindis) mindis=b[i];

	printf("%d \n",mindis);
	
	cudaFree( dev_a ) ;
	cudaFree( dev_b ) ;

	free( a );
   free( b );

	return 0;
}

system · 2013 年5 月 16 日 14:50

我用visual profile运行了一下程序，发现结果如此…不太懂含义

[attach]3132[/attach]

system · 2013 年5 月 16 日 16:07

LZ您好，大致看了您的代码，答复如下：

1：您的代码整体框架是不错的，写的也很清晰和工整，这点很不错。但是在一个warp内规约，且使用shared memory的时候，需要声明shared memory为volatile，否则在计算能力2.x和3.x的卡上会出现错误的结果。如写为 volatile int * s_ptr=sdata;然后使用s_ptr来进行warp内的访存。

这样可以告知编译器每次都立即将值写入shared memory里面，而不是缓冲于寄存器中。
您可以参考CUDA C Programming Guide里面reduction的示例代码。

2：此外，您的总体计算规模偏小，只有10万个数据，计算也不是很复杂，CPU很快就能算完，而GPU还在折腾各种初始化，copy等，所以是不划算的。

3：您的算法在GPU上实现的时候，对于GPU来说，计算较少，整体会卡在访存上，从而您的GPU计算单元不是十分忙碌。

4：您的显卡是GEORCE 605，是一款比较弱的GPU，所以您在对比的时候也需要考虑到这一点。

大致这些，供您参考，祝您编码顺利~

system · 2013 年5 月 16 日 16:13

LZ您好，您的profiler结果大致答复如下：

1：横轴是时间轴，您可以看出来，您的代码的计算时间非常短，而其他开销很大。这一方面是您的计算规模造成的，10万个数据规模偏小。

2：另一方面，您看到前面时间很长的cudaMalloc时间段，其实包含了第一次运行CUDA代码所做的一些环境初始化的工作。建议您再正常的代码前面跑一个空的kernel或者cuda runtime API的函数什么的，使得环境初始化在之前就进行，这俗称“热身”。之后您就可以看出您真正cudaMalloc，cudaMemcpy以及invoke kernel的时间了。这个时候图示才更具指导意义。

大致如此，祝您编码顺利~

system · 2013 年5 月 17 日 12:39

谢谢版主的详细指导啊…我在处理这个问题的时候，发现一个问题是 cudaevent这个计时方式只能用于gpu计算的计时，用于cpu上的计时完全不正确。还有cudaevent的计时单位是毫秒，那么上述的时间上，gpu的计算时间比cpu还是快了那么一点点。

版主您所指出的shared memory的用法上，我已经改正过来，学习了。

还有问题的规模的确是比较小，10w个点的数据，而且是线性的计算量。但是处理这一个点的只是第一步，第二步是要继续处理同样的数据量 1000次（数据不同），正如版主您所说的，主要的问题是卡在传输数据这一步上。我采用了流处理的方法，采用1个流去处理每10w个点，一次4个流，每次传输40w个点到到gpu的内存中，感觉这样的效果不是很明显。

那么就想问版主，有没有好的思路，让传输数据和计算的重叠性比较高？还有请问版主在cpu代码中同时开两个核函数，这两个函数会并行执行吗？

system · 2013 年5 月 17 日 13:14

LZ您好，win7下CPU端计时推荐使用 QueryPerformanceCounter()。

1：“第二步是要继续处理同样的数据量 1000次（数据不同）”这个需要从host传递过来1000次么？还是传递过来一次，在显存里面算1000遍？如果是后者，效率有望提升；如果是前者，那还是一样会卡在传输上的。

2：“我采用了流处理的方法，采用1个流去处理每10w个点，一次4个流…”您采用了4个流来相互掩盖计算和传输么？使用流来掩盖计算和传输是一个好办法，但是如果总的传输时间仍远远大于计算时间的话，效果也不会太好。以及，您能提供您这样使用的大致代码（框架）以及profiler截图么？profiler截图能比较直观地看出您各个stream见重叠的情况如何。

3：“那么就想问版主，有没有好的思路，让传输数据和计算的重叠性比较高？”这个如果算法限制死了，那么只能用多个stream来互相掩盖。以及如前所说，即便stream能完全发挥作用，但是copy时间如果远长于计算时间的话，效果依然不会很好。这是问题算法本身限制的。

4：“ 还有请问版主在cpu代码中同时开两个核函数，这两个函数会并行执行吗？”这个没看明白，可否提供代码示例？

大致如上，祝您编码顺利~

system · 2013 年5 月 17 日 13:35

谢谢版主的回答。QueryPerformanceCounter() 感觉有点高深，试试看。

1、数据是存在文件中的，1000个这样的数据，大概有390M左右，可以一次性传输到显存中，也可以分批次传送过去。我现在采用的办法是分批次传送过去，想要利用传输的时间重叠去计算。但是，390MB的数据量应该很小吧，可能一次性传输会效率更高，我等下试试。

2、我有点不懂，流处理的办法是完全并行的吗？现在的问题是 gpu的计算单元太空闲，而gpu去访问global memory有点太多，瓶颈是这儿，但是流处理能增加gpu的计算单元的计算密度吗？我通过profiler的结果，感觉重合度不是很高，反而感觉就是串行起来了。
[attach]3133[/attach]
4 就是比如
int main()
{
mykernel1<<< >>>( );
mykernel2<<< >>>（）;
}
这两个函数并没有数据相关，但是就是这样顺序下下来。这样能够并行吗？

system · 2013 年5 月 17 日 14:08

LZ您好：

1：390MB的数据确实不多，按照pci-e 2.0 @16X的传输速率计算是390/(500*16)=0.04875s，（仅传输时间，理想状态下，geforce卡可能到不了这个速度）。但是您的kernel计算时间也很短。您可以继续试试看。

2：同一个流内部的运作是串行的，流和流之间是不保证顺序的，从而可以使得计算和copy（指host和device之间的copy）相互重叠。以及还可以让多个运行规模小的kernel之间并行。（不过您的卡比较弱，kernel运行规模不算太小，这里应该kernel并行不了）

stream的“并行”如上解释。

2.1：“现在的问题是 gpu的计算单元太空闲，而gpu去访问global memory有点太多，瓶颈是这儿”这一点和stream无关的，这种情况一般称作“卡在global访存上”，以及这是您的算法本身固有的问题，您的算法是访存密集型的。

2.2：“但是流处理能增加gpu的计算单元的计算密度吗？”不行的，stream机制和您kernel的计算复杂度完全没有关系。

2.3：“我通过profiler的结果，感觉重合度不是很高，反而感觉就是串行起来了。”您的profiler的图看起来太费劲了，请您缩放比例尺，将您kernel附近图示放大观看。以及，请您按照前面“热身”的做法剔除初始化环境的影响。另外大致看到您这只进行了H2D和D2H两次copy，这个和您说的“1000次”似乎不同，以及您的copy并未和计算相重叠，估计是您写法有问题，请提供相关的代码段。

一般来说，假定您是4组“H2D copy——kernel——D2H copy”这样的任务，那么每组指定在一个不同的非默认stream的stream上，使用异步版的copy命令并使用page-locked memory，那么有望使得第二组的H2D copy和第一组的计算时间重合等。以及因为geforce卡不具备双copy引擎，所以具体行为和telsa卡有一定差异。

而您的代码似乎是整体H2D copy——4个stream计算——整体 D2H copy，这样是无法实现计算和copy重合的，这里面有数据依赖性。以及您的图实在难以看清，请适当缩放。

4：不行的，如果您在连续启动这两个kernel的时候是没有设置stream参数的或者设置为同一stream参数，那么他们是串行的，不会自动并行的。如果是在不同的stream上，那么如果kernel运行规模较小，GPU资源充沛的时候，两个kernel可以在GPU上同时运行。（该功能需要fermi或者kepler核心支持）

大致如此，祝您编码顺利~

补充说明了异步版copy实现功能的一个必要条件。

system · 2013 年5 月 17 日 17:17

谢谢版主回答。
发现使用了多次的流处理，效果不见很好。流的传送并没有和计算重叠起来，这是我目前最大的疑问，不知道原因是什么。
附上代码和profiler分析图。

[attach]3134[/attach]

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

const int N =102400; // 数组元素的个数
const int threadsPerBlock = 256;
const int blocksPerGrid = (N+threadsPerBlock-1) / (4*threadsPerBlock ) ;

__global__ void kenel( int2 *a, int *b)	 
{
	extern __shared__ int sdata[];	
	volatile int *s_ptr=sdata;
	
	int tx  = threadIdx.x;	
	int tid = blockIdx.x * blockDim.x + tx;
	int t=blockDim.x*blocksPerGrid;
	
	int minn=1000*1000;
	
	int2 a0;

	a0=a[tid]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;
	
	a0=a[tid+t]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;

	a0=a[tid+2*t]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;

	a0=a[tid+3*t]; 	
	a0.x=a0.x*a0.x+a0.y*a0.y;
	if(minn>a0.x) minn=a0.x;	
	
	s_ptr[tx]=minn;

	__syncthreads();
	
	//if (tx < 256 &&  sdata[tx]> sdata[tx + 256]) sdata[tx]=sdata[tx+256];
	//__syncthreads();
	if (tx < 128 &&  s_ptr[tx]> s_ptr[tx + 128]) s_ptr[tx]=s_ptr[tx+128];
	__syncthreads();
	if (tx <  64 &&  s_ptr[tx]> s_ptr[tx+ 64]) s_ptr[tx]=s_ptr[tx+64];
	__syncthreads();

	if (tx < 32)
	{
			if(s_ptr[tx]> s_ptr[tx + 32]) s_ptr[tx] = s_ptr[tx + 32]; 
			if(s_ptr[tx]> s_ptr[tx + 16]) s_ptr[tx] = s_ptr[tx + 16]; 
			if(s_ptr[tx]> s_ptr[tx + 8])  s_ptr[tx] = s_ptr[tx + 8]; 
			if(s_ptr[tx]> s_ptr[tx + 4])  s_ptr[tx] = s_ptr[tx + 4]; 
			if(s_ptr[tx]> s_ptr[tx + 2])  s_ptr[tx] = s_ptr[tx + 2]; 
			if(s_ptr[tx]> s_ptr[tx + 1])  s_ptr[tx] = s_ptr[tx + 1]; 			
	}	
	b[blockIdx.x]=sdata[0];
	return;
}

const int num=50;
int2 a[num*102400];
int  b[num*blocksPerGrid];
int  ans[500];

//Cpu上空间分配
	//b = (int*)malloc( blocksPerGrid*sizeof(int) );
int main()
{		
	int2 *dev_a;
	int *dev_b;	
	int sharedMemSize=threadsPerBlock*sizeof(int);

	float elapsedTime;    
	cudaEvent_t ss,sp;

	//cpu上产生数据
	for(int j=0;j<num;j++)
		for(int i=0;i<N;i++)
			a[j*N+i].x=rand()%1000,a[j*N+i].y=rand()%1000;
	//cpu上计算最小值
	for(int j=0;j<num;j++)
	{
		int t=1000*1000;
		for(int i=0;i<N;i++)
			if(a[j*N+i].x*a[j*N+i].x+a[j*N+i].y*a[j*N+i].y<t) t=a[j*N+i].x*a[j*N+i].x+a[j*N+i].y*a[j*N+i].y;
		ans[j]=t;
	}
	
	cudaStream_t stream[500];
		
	for(int tt=num;tt<=num;tt+=50)
	{
	
		cudaEventCreate(&ss);
		cudaEventCreate(&sp);
		cudaEventRecord(ss,0);

		//Gpu上空间分配
		cudaMalloc( (void**)&dev_a, tt*N*sizeof(int2) ) ;
		cudaMalloc( (void**)&dev_b, tt*blocksPerGrid*sizeof(int) ) ;
	
		for(int i=0;i<tt;i++)
				cudaStreamCreate(&stream[i]);
  	
		//拷贝输入数据到GPU
		for(int i=0;i<tt;i++)
			cudaMemcpyAsync(dev_a+i*N,a+i*N, N*sizeof(int2),cudaMemcpyHostToDevice,stream[i]);
		//cudaMemcpy(dev_a , a , tt*N*sizeof(int2), cudaMemcpyHostToDevice ) ;	
			
		for(int i=0;i<tt;i++)
			kenel<<<blocksPerGrid,threadsPerBlock,sharedMemSize,stream[i]>>>(dev_a + i* N , dev_b + i * blocksPerGrid );
	
		//kenel<<<blocksPerGrid,threadsPerBlock,sharedMemSize>>>(dev_a  , dev_b );

		//cudaMemcpy(b,dev_b,tt*blocksPerGrid*sizeof(int),cudaMemcpyDeviceToHost);
		for(int i=0;i<tt;i++)
			cudaMemcpyAsync(b+i*blocksPerGrid,dev_b+i*blocksPerGrid,blocksPerGrid*sizeof(int),cudaMemcpyDeviceToHost,stream[i]);
  
		cudaThreadSynchronize();	
		
		//gpu上返回的值，计算最终答案，并比对是否正确
		for(int j=0;j<tt;j++)
		{
			int mindis=1000*1000;
			for(int i=0;i<blocksPerGrid;i++) 
				if(b[j*blocksPerGrid+i]<mindis) mindis=b[j*blocksPerGrid+i];
			if(mindis!=ans[j]) printf("wrong %d %d %d\n",j,ans[j],mindis);
		}

		for(int i=0;i<tt;i++)
			cudaStreamDestroy(stream[i]);

		cudaFree( dev_a ) ;
		cudaFree( dev_b ) ;

		cudaEventRecord(sp,0);
		cudaEventSynchronize(sp);	
		cudaEventElapsedTime(&elapsedTime, ss, sp);
		printf("k=%d Time=%f\n",tt,elapsedTime/1000.0);
		cudaEventDestroy(ss);
		cudaEventDestroy(sp);
	}	
	cudaThreadExit();

	return 0;	
}

system · 2013 年5 月 17 日 17:34

[

LZ您好，手机大致看了下，您主机端没有使用page-locked memory，所以不行。请您修正。

另外cudaThreadSynchronize()是过时的写法，请替换为cudaDeviceSynchronize()

大致如此，欢迎你深夜来访，祝您好运！

system · 2013 年5 月 18 日 09:13

谢谢斑竹。试着用了page-locked memory以后发现的确能够重合了。但是从profiler的分析图来看，为什么中间有这么一段空白时间？是cpu的运行的时间吗？

[attach]3135[/attach]

system · 2013 年5 月 19 日 04:51

LZ您好，因为异步任务是不会立即开始执行的，所以您可以在发布完kernel之后加一个cudaDeviceSynchronize()之类的同步函数以立即让host积攒的命令发布到device端执行。

同时，您每次截图都不甚完整，看起来很费劲，望改进。

最后给您一个参考的链接：
https://developer.nvidia.com/content/how-overlap-data-transfers-cuda-cc

祝您好运！

system · 2013 年5 月 19 日 11:23

谢谢版主。我的profiler图如下，结果显示还是计算密度太小了。我现在觉得拿gpu本身的峰值运算点可能评价写出来的程序合理的方法，而不是与cpu本身的运算时间去比较。但是我不知道怎么去评测gpu本身的峰值运算点，请版主指点下。

[attach]3137[/attach]
[attach]3138[/attach]
[attach]3139[/attach]
[attach]3140[/attach]

system · 2013 年5 月 19 日 16:11

LZ您好，我觉得您这个算法可能卡在访存上了，本身计算密度很小。
以及如此低的compute utilization，使用GPU本身峰值和您的程序比较的话，您能达到的计算峰值依然是很低的。

以及和CPU比较的话，还要考虑不同的CPU，不同的实现和优化效果不同的。

以及，GPU的浮点峰值是SP数量SP频率2。如何达到这个峰值，以及这个峰值是什么含义，请您翻一下旧帖，以前多有讨论的。

祝您编码顺利~

system · 2013 年5 月 20 日 01:28

楼主您可以无视profiler的"Low compute ultilization"提示，

此提示将整个进程的运行时间进行参考，然后用实际的设备占用时间来比较的。

例如您在０时刻进行显存分配，然后中间却没有立刻启动kernel(例如您在host上准备了５秒的数据），然后启动kernel，运算了１秒。只是它会认为您只有20%，提示"low compute ultilization"的。

但如果此时这５秒的host上的准备是无法避免的，那么请无视profiler的此提示。

以及，如果您能反复长期使用设备，那么此提示的使用率会逐步提升。
（例如依然是上例，如果您还是准备了5s数据，但是却反复将您的计算过程重复10000倍(即10000s)。那么您的使用比率将大幅提高。

但这个只要满足要求就好，如果上文说的5s准备无可避免，那么真心可以无视这个的。

感谢您的周末来访。

system · 2013 年5 月 20 日 01:29

以及，这个是对ICE版主观点的补充。

time.h 与cudaEvent记录的时间上差距很大 不知问题所在

大致如此，祝您编码顺利~

time.h 与cudaEvent记录的时间上差距很大不知问题所在