Accelerated Ray Tracing （三）

https://developer.nvidia.com/blog/how-implement-performance-metrics-cuda-cc/

host和device之间是如何同步的。

cudaMemcpy可以进行cudaMemcpyHostToDevice和cudaMemcpyDeviceToHost的同步（阻塞）。数据同步不会被CUDA其他线程打断。但是Kernel运行时异步的。因为数据传输是阻塞的，也就是说如果最后结果从device传输到host，因为数据传输的阻塞性，可以保证之前异步的Kernel已经完成了。
cudaDeviceSynchronize()可以阻塞所有的异步线程直到完成才继续往下执行。 CUDA event API提供了一个简单的计算操作，它包含create、destroy、record和elapse 两个event之间的毫秒。

CUDA事件利用了CUDA streams的概念。CUDA stream 就是在device上按顺序执行的一系列操作。不同stream中的操作可以交错，在某些情况下还可以重叠——这个属性可以用来隐藏host和device之间的数据传输。到目前为止，对GPU的所有操作都是在默认stream，或stream 0(也称为“null stream”)中进行的。

Theoretical Bandwidth

理论带宽可以用产说明中可用的硬件规格来计算。例如，NVIDIA Tesla M2050 GPU使用DDR(双数据率)RAM，具有1546 MHz的内存时钟率和384位宽内存接口。使用这些数据项，NVIDIA Tesla M2050的峰值理论内存带宽为148 GB/秒，如下所示。

在这个计算中，我们转换存储器时钟率到Hz，乘以它的接口宽度(除以8，转换比特到字节)和乘以2由于双数据率。最后，我们除以10^9将结果转换为GB/s。

Effective Bandwidth

我们通过一段时间内特定的程序访问数据量来计算有效带宽。我们使用以下等式。

这里，是以GB/s为单位的有效带宽，R_B是每个内核读取的字节数，W_B是每个内核写入的字节数，t是运行时间，以秒为单位。我们可以修改我们的SAXPY示例来计算有效带宽。

Measuring Computational Throughput

带宽是数据吞吐量的度量。

吞吐量的一个常用度量是GFLOP/s，它代表每秒千兆浮点运算，其中Giga是10^9的前缀。

对于我们的SAXPY计算，度量有效吞吐量很简单:每个SAXPY元素执行一个乘加操作，通常度量为两倍FLOPs，因此我们有

GFLOP/s Effective = 2N / (t * 10^9)

N是SAXPY操作中的元素数，t是经过的时间(以秒为单位)。与理论峰值带宽一样，理论峰值GFLOP/s可以从产品文献中收集。

SAXPY为每个计算的元素读取12个字节，但只执行一个乘法-添加指令(2 FLOPs)，因此很明显它将受到带宽限制，因此在这种情况下(实际上在许多情况下)，带宽是要度量和优化的最重要指标。更常见的做法是使用分析工具来了解计算吞吐量是否成为瓶颈。

#include <iostream>
#include <math.h>
__global__
void saxpy(int n, float a, float* x, float* y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

int main(void)
{
    int N = 20 * (1 << 20);
    float* x, * y, * d_x, * d_y;
    x = (float*)malloc(N * sizeof(float));
    y = (float*)malloc(N * sizeof(float));

    cudaMalloc(&d_x, N * sizeof(float));
    cudaMalloc(&d_y, N * sizeof(float));

    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    cudaEvent_t start, stop;//事件声明
    cudaEventCreate(&start);//事件创建
    cudaEventCreate(&stop);

    cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);//同步
    cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);

    cudaEventRecord(start);//事件记录，也就是start绑定到默认的stream0上，当device到达了stream上的event处，就会记录一个时间戳。

    // Perform SAXPY on 1M elements
    saxpy << <(N + 511) / 512, 512 >> > (N, 2.0f, d_x, d_y);//异步

    cudaEventRecord(stop);

    cudaMemcpy(y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaEventSynchronize(stop);//同步事件,阻塞cpu执行直至stop被记录
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);//计算时间间隔，精度为0.5微秒

    float maxError = 0.0f;
    for (int i = 0; i < N; i++) {
        maxError = max(maxError, abs(y[i] - 4.0f));
    }

    printf("Max error: %fn", maxError);
    printf("Effective Bandwidth (GB/s): %fn", N * 4 * 3 / milliseconds / 1e6);//计算带宽效率我们同时使用理论峰值带宽和观察到的或有效的内存带宽
    //N*4是每个数组读或写所传输的字节数，因子3表示x的读和y的读和写，
    cudaEventDestroy(start);//事件销毁
    cudaEventDestroy(stop);
    cudaFree(x);
    cudaFree(y);
}