add capter5

capter5/ReadMe.md

@@ -0,0 +1,102 @@
+# 获得 GPU 加速的关键
+
+------ 
+
+## CUDA 事件计时
+
+C++ 的计时方法：
+
+1. GCC 和 MSVC 都有的 `clock()`函数；
+2. 原生的 <chrono> 时间库；
+3. GCC 的 `gettimeofday()`计时；
+4. MSVC 的 `QueryPerformanceCounter()` 和 `QueryPerformanceFrequency()` 计时。
+
+CUDA 基于 CUDA 事件的计时方法：
+
+    cudaEvent_t start, stop;
+    CHECK(cudaEventCreate(&start)); // 创建cuda 事件对象。
+    CHECK(cudaEventCreate(&stop));
+    CHECK(cudaEventRecord(start));  // 记录代表开始的事件。
+    cudaEventQuery(start);  // 强制刷新 cuda 执行流。
+
+    // run code.
+
+    CHECK(cudaEventRecord(stop));
+    CHECK(cudaEventSynchronize(stop)); // 强制同步，让主机等待cuda事件执行完毕。
+    float elapsed_time = 0;
+    CHECK(cudaEventElapsedTime(&curr_time, start, stop)); // 计算 start 和stop间的时间差（ms）。
+    printf("host memory malloc and copy: %f ms.\n", curr_time - elapsed_time);  
+
+由于 cuda 程序需要在主机和设备间传递数据，所以当计算强度较小时数据传输的性能对程序总耗时影响更大。
+因此 cuda 的两种浮点数类型对程序性能的影响就较为明显。考虑提供编译选项，指定版本：
+    
+    #ifdef USE_DP
+        typedef double real;  // 双精度
+        const real EPSILON = 1.0e-15;
+    #else
+        typedef float real;   // 单精度
+        const real EPSILON = 1.0e-6f;
+    #endif
+
+在编译时，除了指定 GPU 计算能力 `-arch=sm_50`，还可以指定 c++ 优化等级 `-O3`;同时，可以指定其他  
+编译选项，如 `-DUSE_DP` 启用双精度版本。
+
+    >> nvcc -O3 -arch=sm_50 -DUSE_DP -o ./bin/clock.exe add.cu clock.cu main.cpp
+    ...
+    >> ./bin/clock
+    using double precision version
+    host memory malloc and copy: 2.054112 ms.
+    device memory malloc: 9.063583 ms.
+    kernel function : 0.803360 ms.
+    cuda; no error
+    copy from device to host: 7.489505 ms.  
+
+    >> nvcc -O3 -arch=sm_50 -o ./bin/clock.exe add.cu clock.cu main.cpp
+    ...
+    >> ./bin/clock     
+    host memory malloc and copy: 0.950240 ms.
+    device memory malloc: 5.298208 ms.
+    kernel function : 0.620512 ms.
+    cuda; no errors
+    copy from device to host: 3.034208 ms.
+
+可见双精度版本基本上比单精度版本耗时多一倍。
+
+------
+
+## nvprof 查看程序性能
+
+    >> nvprof ./bin/clock
+
+（没有输出结果）。
+
+------
+
+## 影响 GPU 加速的关键因素
+
+1. 要获得可观的 GPU 加速，就必须尽量缩减主机和设备间数据传输所花时间的占比。
+
+有些计算即使在 GPU 中速度不高也要尽量放在 GPU 中实现，以避免过多数据经由 PCIe 传递。
+
+2. 提高算术强度可以显著地提高 GPU 相对于 CPU 的加速比。
+
+**算术强度**，是指一个计算问题中算术操作的工作量与必要的内存操作的工作量之比。  
+对设备内存的访问速度取决于 GPU 的显存带宽。
+
+3. 核函数的并行规模。
+
+并行规模可以用 GPU 中的线程数目来衡量。  
+一个 GPU 由多个流多处理器SM（streaming multiprocessor）构成，每个 SM 中有若干 CUDA 核心。  
+每个 SM 是相对独立的，一个 SM 中最多驻留的线程数一般为 2048 或 1024（图灵架构）。
+
+若要 GPU 满负荷工作，则核函数中定义的线程总数要不少于某值，一般与 GPU 能够驻留的线程总数相当。
+
+------
+
+## CUDA 的数学函数库
+
+CUDA 提供的数学函数库提供了多种 **数学函数**，同时 CUDA 提供了一些高效率、低准确度的 **内建函数**。
+
+CUDA 数学函数库的更多资料，详见：[CUDA math](https://docs.nvidia.com/cuda/cuda-math-api/index.html).
+
+------

capter5/add.cu

@@ -0,0 +1,51 @@
+#include "add.cuh"
+
+const real c = 3.57;
+
+
+__global__ void add(const real *x, const real *y, real *z, const int N)
+{
+    // 在主机函数中需要依次对每个元素进行操作，需要使用一个循环。
+    // 在设备函数中，因为采用“单指令-多线程”方式，所以可以去掉循环、只要将数组元素索引和线程索引一一对应即可。
+
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    if (n > N) return;
+
+    if (n%5 == 0)
+    {
+        z[n] = add_in_device(x[n], y[n]);
+    }
+    else
+    {
+        add_in_device(x[n], y[n], z[n]);
+    }
+}
+
+__device__ real add_in_device(const real x, const real y)
+{
+    return x + y;
+}
+
+__device__ void add_in_device(const real x, const real y, real &z)
+{
+    z = x + y;
+}
+
+void check(const real *z, const int N)
+{
+    bool has_error = false;
+    for (int i = 0; i < N ;++i)
+    {
+        if (fabs(z[i] - c) > EPSILON)
+        {
+            has_error = true;
+        }
+    }
+
+    printf("cuda; %s\n", has_error ? "has error" : "no error");
+}
+
+
+
+
+

capter5/add.cuh

@@ -0,0 +1,24 @@
+#include "error.cuh"
+
+#include <math.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+#ifdef USE_DP
+    typedef double real;  // 双精度
+    const real EPSILON = 1.0e-15;
+#else
+    typedef float real;   // 单精度
+    const real EPSILON = 1.0e-6f;
+#endif
+
+
+// 核函数。
+__global__ void add(const real *x, const real *y, real *z, const int N);
+
+// 重载设备函数。
+__device__ real add_in_device(const real x, const real y);
+__device__ void add_in_device(const real x, const real y, real &z);
+
+// 主机函数。
+void check(const real *z, const int N);

capter5/clock.cu

@@ -0,0 +1,97 @@
+#include <stdlib.h>
+#include <stdio.h>
+#include "add.cuh"
+#include "clock.cuh"
+
+const real a = 1.23;
+const real b = 2.34;
+
+
+void cuda_clock()
+{
+    const int N = 1e6;
+    const int M = sizeof(real) * N;
+
+    // cuda 计时。
+    float elapsed_time = 0;
+    float curr_time = 0;
+    cudaEvent_t start, stop;
+    CHECK(cudaEventCreate(&start)); // 创建cuda 事件对象。
+    CHECK(cudaEventCreate(&stop));
+    CHECK(cudaEventRecord(start));  // 记录代表开始的事件。
+    cudaEventQuery(start);  // 强制刷新 cuda 执行流。
+
+    // --------------------------------------------------
+    real *h_x, *h_y, *h_z;
+    h_x = new real[N];
+    h_y = new real[N];
+    h_z = new real[N];
+    if (!h_x || !h_y || !h_z)
+    {
+        printf("host memory malloc failed!\n");
+        return;
+    }
+    for (int i = 0; i < N; ++i)
+    {
+        h_x[i] = a;
+        h_y[i] = b;
+    }
+    // --------------------------------------------------
+
+    // 主机申请及初始化内存的耗时。
+    CHECK(cudaEventRecord(stop));
+    CHECK(cudaEventSynchronize(stop)); // 强制同步，让主机等待cuda事件执行完毕。
+    CHECK(cudaEventElapsedTime(&curr_time, start, stop)); // 计算 start 和stop间的时间差（ms）。
+    printf("host memory malloc and copy: %f ms.\n", curr_time - elapsed_time);    
+    elapsed_time = curr_time;
+
+
+    // --------------------------------------------------
+    real *d_x, *d_y, *d_z;
+    CHECK(cudaMalloc(&d_x, M));
+    CHECK(cudaMalloc(&d_y, M));
+    CHECK(cudaMalloc(&d_z, M));
+    CHECK(cudaMemcpy(d_x, h_x, M, cudaMemcpyDefault));
+    CHECK(cudaMemcpy(d_y, h_y, M, cudaMemcpyDefault));
+    // --------------------------------------------------
+
+    // 设备内存申请和拷贝耗时。
+    CHECK(cudaEventRecord(stop));
+    CHECK(cudaEventSynchronize(stop)); 
+    CHECK(cudaEventElapsedTime(&curr_time, start, stop)); 
+    printf("device memory malloc: %f ms.\n", curr_time - elapsed_time);
+    elapsed_time = curr_time;
+
+    // --------------------------------------------------
+    const int block_size = 128;
+    const int grid_size = N/block_size + 1;
+    add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
+    // --------------------------------------------------
+
+    // 核函数运行耗时。
+    CHECK(cudaEventRecord(stop));
+    CHECK(cudaEventSynchronize(stop)); 
+    CHECK(cudaEventElapsedTime(&curr_time, start, stop)); 
+    printf("kernel function : %f ms.\n", curr_time - elapsed_time);
+    elapsed_time = curr_time;   
+
+    // --------------------------------------------------
+    CHECK(cudaGetLastError());
+    CHECK(cudaMemcpy(h_z, d_z, M, cudaMemcpyDefault)); 
+    check(h_z, N);
+    // --------------------------------------------------
+
+    // 数据拷贝耗时。
+    CHECK(cudaEventRecord(stop));
+    CHECK(cudaEventSynchronize(stop)); 
+    CHECK(cudaEventElapsedTime(&curr_time, start, stop)); 
+    printf("copy from device to host: %f ms.\n", curr_time - elapsed_time);
+    elapsed_time = curr_time;  
+
+    if (h_x) delete[] h_x;
+    if (h_y) delete[] h_y;
+    if (h_z) delete[] h_z;
+    CHECK(cudaFree(d_x));
+    CHECK(cudaFree(d_y));
+    CHECK(cudaFree(d_z));
+}

capter5/clock.cuh

@@ -0,0 +1,2 @@
+
+void cuda_clock();

capter5/error.cuh

@@ -0,0 +1,26 @@
+#pragma once
+#include <stdio.h>
+#include <cuda_runtime.h>
+#include <cuda_runtime_api.h>
+#include <device_launch_parameters.h>
+
+#define CHECK(call)                                                     \
+do {                                                                    \
+    const cudaError_t error_code = call;                                \
+    if (error_code != cudaSuccess)                                      \
+    {                                                                   \
+        printf("CUDA ERROR: \n");                                       \
+        printf("    FILE: %s\n", __FILE__);                             \
+        printf("    LINE: %d\n", __LINE__);                             \
+        printf("    ERROR CODE: %d\n", error_code);                     \
+        printf("    ERROR TEXT: %s\n", cudaGetErrorString(error_code)); \
+        exit(1);                                                        \
+    }                                                                   \
+}while(0);                                                              \
+
+
+
+
+
+
+

capter5/main.cpp

@@ -0,0 +1,19 @@
+#include <iostream>
+#include <typeinfo>
+#include "add.cuh"
+#include "clock.cuh"
+
+using namespace std;
+
+
+int main()
+{
+    if (typeid(double) == typeid(real))
+    cout << "using double precision version" << endl;
+
+    cuda_clock();
+
+    return 0;
+}
+
+