【CUDA编程基础】第三章 CUDA程序的基本框架

第零章资料

谭升_CUDA基础_博客

权双_CUDA编程基础入门系列_视频

MAhaitao999_CUDA书籍_pdf 对应的 CUDA编程书本目录

龚大的杂货铺_从上帝视角看GPU_视频

temp:【C++】CUDA期末复习指南上（详细）_cuda c++ 结构体-CSDN博客【C++】CUDA期末复习指南下（详细）_cudasetdevice-CSDN博客

书中源码：GitHub - brucefan1983/CUDA-Programming: Sample codes for my CUDA programming book

CUDA所有版本官网文档：https://developer.nvidia.com/cuda-toolkit-archive

第三章 CUDA程序的基本框架

CUDA程序的基本框架

头文件

常量定义（或宏定义）

C++自定义函数和CUDA核函数的声明（原型）

int main(void)

{

        分配主机与设备内存

        初始化主机中的数据

        将某些数据从主机复制到设备

        调用核函数在设备中计算

        将某些数据从设备复制到主机

        释放主机与设备内存

}

C++自定义函数和CUDA核函数的定义（实现）

`add.cu`

// 第三章
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
/* 不能使用判等确定两个数是否相等，而是使用两数的差值小于某个数。
原因（文心一言）：在 CUDA 程序中，由于浮点数的精度问题，直接使用等号（==）来判断两个浮点数是否相等通常是不准确的。浮点数的表示和计算涉及到舍入误差，这意味着即使两个浮点数在理论上应该是相等的，它们在计算机中的表示可能会略有差异。
需要注意的是，选择合适的容差值是很重要的，它应该根据你的具体应用场景和所需的精度来确定。太小的容差值可能导致误判，而太大的容差值则可能掩盖实际上的差异。
*/
const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.24;
const double c = 3.57;
/*主机函数等同于void add(const double *x, const double *y, double *z, const int N);1、不能同时用__device__和__global__修饰一个函数，即不能将一个函数同时定义为设备函数和核函数。2、也不能同时用__host__和__global__修饰一个函数，即不能将一个函数同时定义为主机函数和核函数。3、编译器决定把设备函数当作内联函数（inline function）或非内联函数，但可以用修饰符__noinline__建议一个设备函数为非内联函数（编译器不一定接受），也可以用修饰符__forceinline__建议一个设备函数为内联函数。
*/
void __host__ add(const double *x, const double *y, double *z, const int N);
// 核函数
void __global__ add_gpu(const double *x, const double *y, double *z);
// 对元素个数与线程块大小不能整除的，进行兼容
void __global__ add_gpu_fit(const double *x, const double *y, double *z, int N);// 设备函数（device function），其不带执行配置（网格数，线程块），只能在设备上执行，只能被核函数或其他设备函数调用
// 版本一：有返回值的设备函数
double __device__ add1_device(const double x, const double y);
// 版本二：用指针的设备函数
void __device__ add2_device(const double x, const double y, double *z);
// 版本三：用引用（reference）的设备函数
void __device__ add3_device(const double x, const double y, double &z);void check(const double *z, const int N);// nvcc chapter3.cu -arch=sm_75 -o chapter3
int main(int argc, char **argv)
{const int N = 100000000;const int M = sizeof(double) * N;double *h_x = (double *)malloc(M);double *h_y = (double *)malloc(M);double *h_z = (double *)malloc(M);double *h_z_cpu = (double *)malloc(M);// 在CPU上执行的主机函数for (int i = 0; i < N; ++i){h_x[i] = a;h_y[i] = b;}add(h_x, h_y, h_z_cpu, N);check(h_z_cpu, N);  int dev = 0;cudaSetDevice(dev);double *d_x, *d_y, *d_z;/* 函数原型cudaError_t cudaMalloc(void **address, size_t size); 其中address是待分配设备内存的指针、size是待分配内存的字节数、cudaError_t是一个错误代号，若如果调用成功，返回cudaSuccess，否则返回一个代表某种错误的代号。*/cudaMalloc((void **)&d_x, M);/*cudaMalloc() 函数使用双重指针作为变量是因为该函数的功能是改变指针d_x 本身的值（将一个指针赋值给d_x），而不是改变d_x 所指内存缓冲区中的变量值。在这种情况下，必须将d_x 的地址&d_x 传给函数cudaMalloc() 才能达到此效果。从另一个角度来说，函数cudaMalloc() 要求用传双重指针的方式改变一个指针的值，而不是直接返回一个指针，是因为该函数已经将返回值用于返回错误代号，而C++ 又不支持多个返回值。*/cudaMalloc((void **)&d_y, M);cudaMalloc((void **)&d_z, M);cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);const int block_size = 128;// 可以被整除，网格数为整除，若不能，则需多申请一个网格，需做兼容const int grid_size = N / block_size;const int grid_size_fit = (N - 1) / block_size + 1;printf("block_size:%d grid_size:%d \n", block_size, grid_size);printf("block_size:%d grid_size_fit:%d \n", block_size, grid_size);// block_size:128 grid_size:781250add_gpu<<<grid_size, block_size>>>(d_x, d_y, d_z);add_gpu_fit<<<grid_size_fit, block_size>>>(d_x, d_y, d_z, N);/*cudaError_t cudaMemcpy(void *dst, const void *src,size_t count,enum cudaMemcpyKind kind);dst是目标地址、src是源地址、count是复制数据的字节数、kind一个枚举类型的变量，标志数据传递方向有以下几种cudaMemcpyHostToHost、cudaMemcpyHostToDevice、cudaMemcpyDeviceToHost、cudaMemcpyDeviceToDevice分别表示从主机到主机、从主机到设备、从设备到主机、从设备到设备，还有cudaMemcpyDefault，表示根据指针dst和src所指地址自动判断数据传输的方向。这要求系统具有统一虚拟寻址（unified virtual addressing）的功能（要求64位的主机）cudaError_t是一个错误代号，若如果调用成功，返回cudaSuccess，否则返回一个代表某种错误的代号。*/cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);check(h_z, N);// 释放内存free(h_x);free(h_y);free(h_z);// 释放设备内存（显存）cudaFree(d_x);cudaFree(d_y);cudaFree(d_z); return 0;
}
/*将主机中的函数改为设备中的核函数：基本上就是去掉一层循环。在设备的核函数中，我们用“单指令-多线程”的方式编写代码，故可去掉该循环，只需将数组元素指标与线程指标一一对应即可。
*/
void __global__ add_gpu(const double *x, const double *y, double *z)
{// 全局线程idconst int n = blockDim.x * blockIdx.x + threadIdx.x;// printf("thread idx: %d\n",n);// 数组元素指标与线程指标一一对应z[n] = x[n] + y[n];
}void __global__ add_gpu_fit(const double *x, const double *y, double *z, int N)
{// 全局线程idconst int n = blockDim.x * blockIdx.x + threadIdx.x;// printf("thread idx: %d\n",n);// 核函数虽不允许有返回值，但可以使用return语句提前结束核函数执行。if (n >= N)return;z[n] = x[n] + y[n];
}double __device__ add1_device(const double x, const double y)
{return (x + y);
}void __device__ add2_device(const double x, const double y, double *z)
{*z = x + y;
}void __device__ add3_device(const double x, const double y, double &z)
{z = x + y;
}void add(const double *x, const double *y, double *z, const int N)
{for (int i = 0; i < N; ++i){z[i] = x[i] + y[i];}
}void check(const double *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("%s\n", has_error ? "Has errors" : "No errors");
}

函数名无特殊要求，而且支持C++中的重载（overload），即可以用同一个函数名表示具有不同参数列表的函数。
除非使用统一内存编程机制（将在第12章介绍），否则传给核函数的数组（指针）必须指向设备内存。
核函数不可成为一个类的成员。通常的做法是用一个包装函数调用核函数，而将包装函数定义为类的成员。
无论是从主机调用，还是从设备调用，核函数都是在设备中执行。调用核函数时必须指定执行配置，即三括号和它里面的参数。
从计算能力3.5开始，引入动态并行（dynamic parallelism）机制，在核函数内部可以调用其他核函数，以及进行递归。

拓展——指针

在C++中，当你声明一个指针时，你实际上正在创建一个新的变量，这个变量能够存储内存地址。具体来说，以下是当你声明一个指针时所发生的事情：

        分配内存：系统为指针变量本身分配内存。这个内存的大小取决于系统的位数（例如，在32位系统上，指针通常占用4个字节；在64位系统上，指针通常占用8个字节）。请注意，这里分配的内存是用于存储地址的，而不是用于存储指针所指向的数据。

        确定类型：在声明指针时，你必须指定指针的类型（例如，int*、double*、char* 等）。这告诉编译器该指针可以指向哪种类型的数据。指针的类型对于后续的内存访问和解引用操作至关重要，因为它决定了如何解释存储在指针所指向地址上的数据。

        初始化：你可以选择初始化指针，即给它赋予一个具体的内存地址。如果你没有立即初始化指针，那么它的初始值是未定义的（也就是说，它可能包含任何值），这通常被称为“野指针”。在使用指针之前，最好总是将其初始化为nullptr（表示“没有指向任何对象”）或一个有效的内存地址。

例如，以下是一个简单的C++指针声明：

        int *ptr = nullptr; // 声明一个指向整数的指针，并将其初始化为nullptr

        在这个例子中，ptr是一个能够存储整数类型数据内存地址的指针变量。它被初始化为nullptr，表示它目前不指向任何有效的内存位置。后续你可以通过赋值操作使其指向一个具体的整数变量或动态分配的内存块。

-----------------------------------------------------------------------------------------------------------------

双重指针（也称为二级指针或指向指针的指针）在C++中的声明和使用与单重指针有相似之处，但也有其特殊之处。以下是关于双重指针声明时的详细解释：

        1. 内存分配：与单重指针一样，当声明一个双重指针时，系统也会为其分配内存。这块内存用于存储另一个指针的地址，而不是直接存储数据。因此，双重指针本身占用的内存大小仍然取决于系统的位数。

        2. 类型确定：双重指针的类型指明了它所指向的对象类型，即另一个指针的类型。例如，int** 表示一个指向int*类型指针的指针。这告诉编译器，通过双重指针解引用后得到的是一个指向整数的指针。

        3. 初始化：双重指针同样需要初始化，以避免成为野指针。它可以被初始化为nullptr，表示它不指向任何指针，或者初始化为另一个指针变量的地址。未初始化的双重指针可能导致未定义的行为，甚至程序崩溃。

        4. 用途：双重指针在编程中有多种用途。它们常用于动态内存分配（如二维数组的动态分配）、修改指针本身的值（如在函数内部改变外部指针的指向），以及构建复杂的数据结构（如树、图等）。

        5. 操作与解引用：使用双重指针时，需要进行两次解引用操作才能访问到最终的数据。第一次解引用得到的是指向数据的指针，第二次解引用才能得到数据本身。因此，操作双重指针时需要格外注意指针的层级和指向。

综上所述，双重指针在声明时与单重指针类似，都需要分配内存、确定类型，并最好进行初始化。然而，双重指针的用途和操作方式更为复杂，需要程序员有更深入的理解和谨慎的操作。