From 7e6d50d4735470c6d5c7da84f5c83dd41eefbef3 Mon Sep 17 00:00:00 2001 From: unknown Date: Mon, 15 Nov 2021 13:04:38 +0800 Subject: [PATCH] add capter3 --- capter3/ReadMe.md | 38 ++++++++++++++++---- capter3/add.cu | 89 ++++++++++++++++++++++++++++++++++++++--------- 2 files changed, 105 insertions(+), 22 deletions(-) diff --git a/capter3/ReadMe.md b/capter3/ReadMe.md index b7b17df..5bd48be 100644 --- a/capter3/ReadMe.md +++ b/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__` 建议编译器将一个设备函数当作内联函数。 + +设备函数可以有返回值。 + +------ diff --git a/capter3/add.cu b/capter3/add.cu index 0bae96a..e964c9a 100644 --- a/capter3/add.cu +++ b/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<<>>(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"); }