CUDA实战笔记(10)——使用事件测量性能

本系列代码托管于：https://github.com/chintsan-code/cuda_by_example

本篇使用的项目为：ray_tracing_noconst_event、ray_tracing_const_event

一. 常量内存带来的性能提升

__constant__将把变量的访问限制只读，与全局内存相比，从常量内存中读取相同的数据可以节省内存带宽，这是因为：

• 对常量内存的单次读操作可以广播到其他的”邻近(Nearby)”的Thread，这将节约15次读取操作（为什么是15，在下面有说）
• 常量内存的数据将缓存起来，因此对相同地址的连续读操作将不会产生额外的内存通信量。

如何理解”邻近(Nearby)”？首先要引入一个新的概念：线程束(Warp)。Warp可以看做是一组Thread通过交织形成一个整体。在CUDA架构中，Wrap是指一个包含32个Thread的集合，这个线程几个被”编织在一起”并且以”步调一致(Lockstep)”的形式执行。在程序中的每一行，Warp中的每个Thread都将在不同的数据上执行相同的指令。

当处理常量内存时，NVIDA硬件将把单次内存读取操作广播到每个半线程束(Half-Warp)。在Half-Warp中包含16个Thread，即Warp的一半。如果在Half-Warp中每个Thread都从常量内存的相同地址读取数据，那么GPU只会产生一次读取请求并在随后将读取到的数据广播到每个Thread。如果从常量内存中读取大量的数据，那么这种方法产生的内存流量只是使用全局内存时的1/16（大约6%）

在读取常量内存的数据时，所节约的并不仅限于减少94%的带宽，由于这块内存的中的数据时不会发生变化的，因此硬件将主动把这个常量数据缓存在GPU上。在第一次从常量内存的某个地址上读取后，当其他Half-Warp请求同一地址时，那么将会命中缓存，这也会减少内存流量。

在缓存数据之后，其他线程将不会产生内存流量，原因有2个：

• 线程将在Half-Warp的广播中收到这个数据。
• 从常量内存缓存中收到数据

需要注意的是，Half-Warp是一把双刃剑，当所有16个Thread都读取相同地址的数据时，可以极大地提升性能，但是当所有16个Thread分别读取不同地址的数据时，它实际上会降低性能。

只有当16个线程都需要相同的读取请求时，才值得使用Half-Warp，将读取操作广播到16个Thread。如果这16个Thread需要访问常量内存中的不同数据，那么这16次读取操作就会被串行化，从而需要16倍的时间来发出请求。但如果从全局内存中读取，那么这些请求会同时发出，二者比较之下，反而是使用全局内存中读取更快。

二. 使用事件来测量性能

测量性能的常见手段是测量程序的执行时间，看看哪个版本的执行时间更短。我们可以使用CPU或操作系统提供的某个计时器，但是这将带来各种延时(包括操作系统线程调度，高精度CPU计时器可用性等方面)。而且，当GPU核函数运行时，我们还可以在主机上异步地执行计算。测量这些主机运算时间的唯一方式是使用CPU或操作系统的定时机制。为了测量GPU在某个任务上花费的时间，我们将使用CUDA的事件API。

CUDA中的事件本质上是一个GPU时间戳，这个时间戳是在用户指定的时间上记录的。由于GPU本身支持记录时间戳，因此避免了当使用CPU定时器来统计GPU执行的时间时可能遇到的诸多问题。

获取一个GPU时间戳只需要2个步骤：

• 首先创建一个事件
• 然后记录一个事件

代码如下：

cudaEvent_t start;
cudaEventCreate(&start);  // 创建一个事件
cudaEventRecord(start, 0);  // 记录一个事件

注意到cudaEventRecord中还指定了第二个参数0，当前不需要管，后面流(Stream)部分继续。

要记录一段时间，不仅要创建起始时间start，还要创建结束事件stop：

cudaEvent_t start, stop;
cudaEventCreate(&start);  // 创建一个事件
cudaEventCreate(&stop);
cudaEventRecord(start, 0);  // 记录一个事件

// 在GPU上执行了一些工作

cudaEventRecord(stop, 0);

然而，这样依然会存在一个问题：当我们在CUDA C中执行某些异步函数时，GPU会执行相关代码，但同时，CPU也会执行。

因此需要有一种方式告诉CPU同步：

cudaEvent_t start, stop;
cudaEventCreate(&start);  // 创建一个事件
cudaEventCreate(&stop);
cudaEventRecord(start, 0);  // 记录一个事件

// 在GPU上执行了一些工作

cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);  // 在事件stop同步

现在，程序告诉运行时要阻塞后面的语句，知道GPU执行到达stop事件。当cudaEventSynchronize返回时，就表示在stop事件之前的所有GPU工作都已经完成了，因此可以安全地读取在stop中保存的时间戳。

需要注意的是，由于CUDA事件是直接在GPU上实现的，因此它们不适用与对同时包含设备代码和主机代码的混合代码计时，也就是说，如果试图通过CUDA事件对核函数和设备内存复制之外的代码进行计时，将得到不可靠的结果。

之后，需用计算两个事件之间经历的时间：

float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
printf("Time to generate:  %3.1f ms\n", elapsedTime);

最后需要手动销毁之前创建的事件，类似使用free释放掉使用malloc分配的内存，每个事件都需要进行一次销毁：

cudaEventDestroy(start);
cudaEventDestroy(stop);

完整代码如下：

// ray_tracing_const_event

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include <stdio.h>
#include <math.h>
#include "../../common/book.h"
#include "../../common/cpu_bitmap.h"

#define INF 2e10f
#define rnd( x ) (x * rand() / RAND_MAX)
#define DIM 1024
#define SPHERES 20

struct Sphere
{
    float r, g, b;
    float radius;
    float x, y, z;

    __device__ float hit(float ox, float oy, float* n) {
        float dx = ox - x;
        float dy = oy - y;
        if (dx * dx + dy * dy < radius * radius) {
            float dz = sqrtf(radius * radius - dx * dx - dy * dy);
            *n = dz / sqrtf(radius * radius);
            return dz + z;
        }
        return -INF;
    }
};

// 常量内存，需要静态分配
__constant__ Sphere dev_s[SPHERES];

__global__ void kernel(unsigned char* ptr) {
    int x = blockIdx.x * blockDim.x + threadIdx.x;
    int y = blockIdx.y * blockDim.y + threadIdx.y;
    int offset = y * gridDim.x * blockDim.x + x;

    float ox = (x - DIM / 2);
    float oy = (y - DIM / 2);

    float r = 0, g = 0, b = 0;
    float maxz = -INF;
    for (int i = 0; i < SPHERES; i++) {
        float n;
        float t = dev_s[i].hit(ox, oy, &n);
        if (t > maxz) {
            float fscale = n;
            r = dev_s[i].r * fscale;
            g = dev_s[i].g * fscale;
            b = dev_s[i].b * fscale;
            maxz = t;
        }
    }

    ptr[offset * 4 + 0] = (int)(r * 255);
    ptr[offset * 4 + 1] = (int)(g * 255);
    ptr[offset * 4 + 2] = (int)(b * 255);
    ptr[offset * 4 + 3] = 255;
}


int main() {
    CPUBitmap bitmap(DIM, DIM);
    unsigned char* dev_ptr;

    // 记录起始时间
    cudaEvent_t start, stop;
    HANDLE_ERROR(cudaEventCreate(&start));  // 创建一个事件
    HANDLE_ERROR(cudaEventCreate(&stop));
    HANDLE_ERROR(cudaEventRecord(start, 0));  // 记录一个事件


    HANDLE_ERROR(cudaMalloc((void**)&dev_ptr, bitmap.image_size()));

    Sphere* spheres = (Sphere*)malloc(SPHERES * sizeof(Sphere));
    for (int i = 0; i < SPHERES; i++) {
        spheres[i].r = rnd(1.0f);
        spheres[i].g = rnd(1.0f);
        spheres[i].b = rnd(1.0f);
        spheres[i].x = rnd(1000.0f) - 500;
        spheres[i].y = rnd(1000.0f) - 500;
        spheres[i].z = rnd(1000.0f) - 500;
        spheres[i].radius = rnd(100.0f) + 20;
    }
    HANDLE_ERROR(cudaMemcpyToSymbol(dev_s, spheres, SPHERES * sizeof(Sphere)));
    free(spheres);

    dim3 blocks(DIM / 16, DIM / 16);
    dim3 threads(16, 16);
    kernel<<<blocks, threads>>>(dev_ptr);

    HANDLE_ERROR(cudaMemcpy(bitmap.get_ptr(), dev_ptr,
        bitmap.image_size(),
        cudaMemcpyDeviceToHost));

    // 获取结束时间，并显示计时结果
    HANDLE_ERROR(cudaEventRecord(stop, 0));  // 记录事件
    HANDLE_ERROR(cudaEventSynchronize(stop));  // 阻塞后面的语句，直到GPU执行到达stop事件
    
    float elapsedTime;
    HANDLE_ERROR(cudaEventElapsedTime(&elapsedTime, start, stop));
    printf("Time to generate:  %3.1f ms\n", elapsedTime);

    // 销毁事件，类似对malloc()分配的内存进行free()
    HANDLE_ERROR(cudaEventDestroy(start));
    HANDLE_ERROR(cudaEventDestroy(stop));

    bitmap.display_and_exit();

    HANDLE_ERROR(cudaFree(dev_ptr));

    return 0;
}

---

参考：

• 《GPU高性能编程 CUDA实战》