6.3_使用事件来测量性能

6.3 使用事件来测量性能

在充分理解了常量内存既可能带来正面影响，也可能带来负面影响后，你已经决定将光线跟踪器修改为使用常量内存。但是，如何判断常量内存对程序性能有着多大的影响？最简单的衡量标准之一就是回答这个问题：哪个版本的执行时间更短？我们可以使用CPU或者操作系统中的某个计时器，但这将带来各种延迟（包括操作系统线程调度，高精度CPU计时器可用性等方面）。而且，当GPU核函数运行时，我们还可以在主机上异步地执行计算。测量这些主机运

算时间的唯一方式是使用CPU或者操作系统的定时机制。为了测量GPU在某个任务上花费的时间，我们将使用CUDA的事件API。

CUDA中的事件本质上是一个GPU时间戳，这个时间戳是在用户指定的时间点上记录的。由于GPU本身支持记录时间戳，因此就避免了当使用CPU定时器来统计GPU执行的时间时可能遇到的诸多问题。这个API使用起来很容易，因为获得一个时间戳只需要两个步骤：首先创建一个事件，然后记录一个事件。例如，在某段代码的开头，我们告诉CUDA运行时记录当前时间。首先创建一个事件，然后记录这个事件：

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

你将注意到，当告诉运行时记录事件start时还指定了第二个参数。在前面的示例中，这个参数为0。这个参数当前并不重要，因此在这里不进行解释。如果你确实好奇，那么可以在我们讨论流（Stream）时再介绍这个参数。

要统计一段代码的执行时间，不仅要创建一个起始事件，还要创建一个结束事件。当在GPU上执行某个工作时，我们不仅要告诉CUDA运行时记录起始时间，还要记录结束时间：

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

// 在GPU上执行一些工作

cudaEventRecord( stop, 0 );

然而，当按照这种方式来统计GPU代码的执行时间时，仍然存在一个问题。要修复这个问题，只需一行代码，但需要对这行代码进行一些解释。在使用事件时，最复杂的情况是当我们在CUDA C中执行的某些异步函数调用时。例如，当启动光线跟踪器的核函数时，GPU开始执行代码，但在GPU执行完之前，CPU会继续执行程序中的下一行代码。从性能的角度来看，这是非常好的，因为这意味着我们可以在GPU和CPU上同时进行计算，但从逻辑概念上来看，这将使计时工作变得更加复杂。

你应该将cusdaEventRecord()视为一条记录当前时间的语句，并且把这条语句放入GPU的未完成工作队列中。因此，直到GPU执行完了在调用cusdaEventRecord()之前的所有语句时，事件才会被记录下来。由stop事件来测量正确的时间正是我们所希望的，但仅当GPU完成了之前的工作并且记录了stop事件后，才能安全地读取stop时间值。幸运的是，我们有一种方式告诉CPU在某个事件上同步，这个事件API函数就是cusdaEventSynchronize():

cudAAEvent_t start,stop;   
cudAAEventCreate(&start);   
cudAAEventCreate(&stop);   
cudAAEventRecord( start,0）;

//在GPU上执行一些工作

cudaEventRecord( stop, 0 );  
cudaEventSynchronize( stop );

现在，我们已经告诉运行时阻塞后面的语句，直到GPU执行到达stop事件。当CUDAEventSynchronize返回时，我们知道在stop事件之前的所有GPU工作已经完成了，因此可以安全地读取在stop中保存的时间戳。值得注意的是，由于CUDA事件是直接在GPU上实现的，因此它们不适用于对同时包含设备代码和主机代码的混合代码计时。也就是说，如果你试图通过CUDA事件对核函数和设备内存复制之外的代码进行计时，将得到不可靠的结果。

测量光线跟踪器的性能

为了对光线跟踪器计时，我们需要分别创建一个起始事件和一个结束事件。下面是一个支持计时的光线跟踪器，其中没有使用常量内存：

int main(void) { //记录起始时间 CUDAEvent_t start, stop; HANDLE_ERROR( CUDAEventCreate( &start)); HANDLE_ERROR( CUDAEventCreate( &stop ); HANDLE_ERROR( CUDAEventRecord( start, 0 ) ); CPUBitmap bitmap(DIM, DIM); unsigned char *dev_bitmap; //在GPU上为输出位图分配内存 HANDLE_ERROR( CUDA malloc( (void**)&dev_bitmap, bitmap(image_size()) ); //为Sphere数据集分配内存 HANDLE_ERROR( CUDA malloc( (void**)&s, sizeof(Sphere) * SPHERES ) ); //分配临时内存，对其初始化，复制到 //GPU上的内存，然后释放临时内存 Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES ); for (int i=0; i<SPHERES; i++) { temp_s[i].r = rnd(1.0f); temp_s[i].g = rnd(1.0f);
temp_s[i].b = rnd(1.0f);   
temp_s[i].x = rnd(1000.0f) - 500;   
temp_s[i].y = rnd(1000.0f) - 500;   
temp_s[i].z = rnd(1000.0f) - 500;   
temp_s[i].radius = rnd(100.0f) + 20;   
}   
HANDLE_ERROR(udaMemcpy(s, temp_s, sizeof(Sphere) * SPHERES,udaMemcpyHostToDevice));   
free(temp_s);

//从球面数据中生成一张位图

dim3 grids(DIM/16, DIM/16);  
dim3 threads(16, 16);  
kernel<<<grids, threads>>>(s, dev bitmap);

//将位图从GPU上复制回来并显示

HANDLE_ERROR(udaMemcpy( bitmap.get_ptr(),dev bitmap, bitmap,image_size(),udaMemcpyDeviceToHost)）;

//获得结束时间，并显示计时结果

HANDLE_ERROR(udaEventRecord(stop,0)）;   
HANDLE_ERROR(udaEventSynchronize(stop))；   
float elapsedTime;   
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime, start，stop）);   
printf("Time to generate: %3.1f ms\n"， elapsedTime）;   
HANDLE_ERROR(udaEventDestroy（start）);   
HANDLE_ERROR(udaEventDestroy(stop）);

//显示位图

bitsmap.display_and_exit();

//释放内存

cudaFree( dev bitmap );  
cudaFree(s);

注意，我们调用了两个额外的函数，分别为cusdaEventElapsedTime()和cusdaEventDestroy()。cusdaEventElapsedTime()是一个工具函数，用来计算两个事件之间经历的时间。第一个参数为某个浮点变量的地址，在这个参数中将包含两次事件之间经历的时间，单位为毫秒。

当使用完事件后，需要调用)cudaEventDestroy()来销毁它们。这相当于对malloc()分配的内存调用free()，因此每个cudaEventCreate()都对应一个cudaEventDestroy()同样是非常重要的。

接下来对使用常量内存的光线跟踪器进行计时：

int main(void) { //记录起始时间 CUDAEvent_t start, stop; HANDLE_ERROR( CUDAEventCreate( &start ) ); HANDLE_ERROR( CUDAEventCreate( &stop ) ); HANDLE_ERROR( CUDAEventRecord( start, 0 ) ); CPUBitmap bitmap(DIM, DIM); unsigned char *dev_bitmap; //在GPU上为输出位图分配内存 HANDLE_ERROR( CUDA malloc( (void**)&dev_bitmap, bitmap(image_size()) ); //分配临时内存，对其初始化、复制到 //GPU上的内存，然后释放临时内存 Sphere *temp_s = (Sphere*)malloc( sizeof(Sphere) * SPHERES); for (int i=0; i<SPHERES; i++) { temp_s[i].r = rnd(1.0f); temp_s[i].g = rnd(1.0f); temp_s[i].b = rnd(1.0f); temp_s[i].x = rnd(1000.0f) - 500; temp_s[i].y = rnd(1000.0f) - 500; temp_s[i].z = rnd(1000.0f) - 500; temp_s[i].radius = rnd(100.0f) + 20; } HANDLE_ERROR( CUDAMemcpyToSymbol(s, temp_s, sizeof(Sphere) * SPHERES)); free(temp_s);

//从球面数据中生成一张位图

dim3 grids(DIM/16, DIM/16);  
dim3 threads(16, 16);  
kernel<<<grids, threads>>>(dev bitmap);

// 将位图从GPU上复制回来并显示

HANDLE_ERROR(udaMemcpy(map.get_ptr(),dev bitmap, bitmap,image_size(),udaMemcpyDeviceToHost）);
//获得结束时间，并显示计时结果
HANDLE_ERROR(udaEventRecord(stop,0) );
HANDLE_ERROR(udaEventSynchronize(stop) );
float elapsedTime;
HANDLE_ERROR(udaEventElapsedTime(& elapsedTime,
start,stop) );
printf("Time to generate: %3.1f ms\n", elapsedTime);
HANDLE_ERROR(udaEventDestroy(start) );
HANDLE_ERROR(udaEventDestroy(stop) );
//显示
 bitmap.display_and_exit(   );
//释放内存
CUDAFree( dev嵋 );
}

现在，当我们运行这两个版本的光线跟踪器时，就可以比较在GPU上完成相同工作时节约的时间。这将告诉我们，引入常量内存是提升应用程序的性能还是降低性能。在这里的情况中，使用常量内存将极大地提升性能。我们在GeForce GTX 280上进行了实验，实验结果表明，使用常量内存的光线跟踪器的性能比使用全局内存的性能提升了 50%50\% 。在不同的GPU上，你得到的实验结果可能会有所不同，但使用常量内存的光线跟踪器应该比不使用常量内存的光线跟踪器要更快一些。

Previous6.2_常量内存
Next6.4_本章小结