add capter3

capter3/ReadMe.md

@@ -14,19 +14,45 @@
 
     int main()
     {
-        分配主机和设备内存
-        初始化主机中数据
-        将某些数据从主机复制到设备
-        调用核函数在设备中计算
-        将某些数据从设备复制到主机
-        释放主机和设备内存
+        1. 分配主机和设备内存
+        2. 初始化主机中数据
+        3. 将某些数据从主机复制到设备
+        4. 调用核函数在设备中计算
+        5. 将某些数据从设备复制到主机
+        6. 释放主机和设备内存
     }
 
     c++ 自定义函数和 cuda 核函数的定义
 
+CUDA 核函数的要求：
 
+1. 返回类型必须是 `void`，但是函数中可以使用 `return`（但不可以返回任何值）；
+2. 必须使用限定符 `__glolbal__`，也可以加上 c++ 限定符；
+3. 核函数支持 c++ 的重载机制；
+4. 核函数不支持可变数量的参数列表，即参数个数必须确定；
+5. 一般情况下，传给核函数的数组（指针）必须指向设备内存（“统一内存编程机制”除外）；
+6. 核函数不可成为一个类的成员（一般以包装函数调用核函数，将包装函数定义为类成员）；
+7. 在计算能力3.5之前，核函数之间不能相互调用；之后，通过“动态并行”机制可以调用；
+8. 无论从主机调用还是从设备调用，核函数都在设备中执行（“<<<,>>>”指定执行配置）。
 
+------
 
+## 自定义设备函数
 
+核函数可以调用不带执行配置的自定义函数，即 **设备函数**。
 
+设备函数在设备中执行、在设备中被调用；而核函数在设备中执行、在主机中被调用。
 
+1. `__global__`修饰的函数称为核函数，一般由主机调用、在设备中执行；
+2. `__device__`修饰的函数称为设备函数，只能被核函数或其他设备函数调用、在设备中执行；
+3. `__host__`修饰主机段的普通 c++ 函数，在主机中被调用、在主机中执行，一般可以省略；
+4. 可以同时用 `__host__` 和 `__device__` 修饰函数，从而减少代码冗余，此时编译器将
+分别在主机和设备上编译该函数；
+5. 不能同时用 `__global__` 和 `__device__` 修饰函数；
+6. 不能同时用 `__global__` 和 `__host__` 修饰函数；
+7. 可以通过 `__noinline__` 建议编译器不要将一个设备函数当作内联函数；
+8. 可以通过 `__forceinline__` 建议编译器将一个设备函数当作内联函数。
+
+设备函数可以有返回值。
+
+------

capter3/add.cu

@@ -9,7 +9,14 @@ const double b = 2.34;
 const double c = 3.57;
 
 
+// 核函数。
 __global__ void add(const double *x, const double *y, double *z, const int N);
+
+// 重载设备函数。
+__device__ double add_in_device(const double x, const double y);
+__device__ void add_in_device(const double x, const double y, double &z);
+
+// 主机函数。
 void check(const double *z, const int N);
 
 
@@ -18,35 +25,84 @@ int main()
     const int N = 1e4;
     const int M = sizeof(double) * N;
 
-    // 申请内存。
-    double *x = (double*) malloc(M);
-    double *y = (double*) malloc(M);
-    double *z = (double*) malloc(M);
+    // 申请主机内存。
+    // 支持使用 new-delete 方式创建和释放内存。
+    //double *h_x = (double*) malloc(M);
+    double *h_x = new double[N];
+    double *h_y = (double*) malloc(M);
+    double *h_z = (double*) malloc(M);
 
+    // 初始化主机数据。
     for (int i = 0; i < N; ++i)
     {
-        x[i] = a;
-        y[i] = b;
+        h_x[i] = a;
+        h_y[i] = b;
     }
 
-    add(x, y, z, N);
-    check(z, N);
+    // 申请设备内存。
+    // cudeError_t cudaMalloc(void **address, size_t size);
+    double *d_x, *d_y, *d_z;
+    cudaMalloc((void**)&d_x, M);
+    cudaMalloc((void**)&d_y, M);
+    cudaMalloc((void**)&d_z, M);
 
-    // 释放内存。
-    free(x);
-    free(y);
-    free(z);
+    // 从主机复制数据到设备。
+    // cudaError_t cudaMemcpy(void *dst, void *src, size_t count, enum cudaMemcpyKind kind);
+    // kind 可以简化使用 `cudaMemcpyDefault`，由系统自动判断拷贝方向（x64主机）。
+    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
+    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
+
+    // 在设备中执行计算。
+    const int block_size = 128;
+    const int grid_size = N/128 + 1; // 线程数应该不少于计算数目。
+    add<<<grid_size, block_size>>>(d_x, d_y, d_z, N);
+
+    // 从设备复制数据到主机。
+    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
+    check(h_z, N);
+
+    // 释放主机内存。
+    // free(h_x);
+    if (h_x) delete[] h_x;
+    free(h_y);
+    free(h_z);
+
+    // 释放设备内存。
+    // cudaError_t cudaFree(void *address);
+    cudaFree(d_x);
+    cudaFree(d_y);
+    cudaFree(d_z);
 
     return 0;
 }
 
 
-void add(const double *x, const double *y, double *z, const int N)
+__global__ void add(const double *x, const double *y, double *z, const int N)
 {
-    for (int i = 0; i < N; ++i)
+    // 在主机函数中需要依次对每个元素进行操作，需要使用一个循环。
+    // 在设备函数中，因为采用“单指令-多线程”方式，所以可以去掉循环、只要将数组元素索引和线程索引一一对应即可。
+
+    const int n = blockDim.x * blockIdx.x + threadIdx.x;
+    if (n > N) return;
+
+    if (n%5 == 0)
     {
-        z[i] = x[i] + y[i];
+        z[n] = add_in_device(x[n], y[n]);
     }
+    else
+    {
+        add_in_device(x[n], y[n], z[n]);
+    }
+}
+
+__device__ double add_in_device(const double x, const double y)
+{
+    return x + y;
+}
+
+__device__ void add_in_device(const double x, const double y, double &z)
+{
+    z = x + y;
 }
 
 void check(const double *z, const int N)
@@ -56,9 +112,10 @@ void check(const double *z, const int N)
     {
         if (fabs(z[i] - c) > EPSILON)
         {
+            //printf("%d, %f, %f\n", i, z[i], c);
             has_error = true;
         }
     }
 
-    printf("%s\n", has_error ? "has error" : "no error");
+    printf("cuda; %s\n", has_error ? "has error" : "no error");
 }