CUDA编程（三）CUDA程序的基本框架

1.CUDA程序基本框架

头文件包含

常量定义（或者宏定义）

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


int main(void) {

分配主机与设备内存

初始化主机中的数据

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

调用核函数在设备中进行计算

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

释放主机与设备内存

}

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

数据量越大，网格越大。在linux中使用nvcc编译程序时，需要指定计算能力，才能在合适的网格最大值范围。

在CUDA运行时API中，没有明显地初始化设备（即GPU）的函数。在第一次调用一 个和设备管理及版本查询功能无关的运行时API函数时，设备将自动地初始化。

2.例子，两数相加

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include<math.h>
#include<stdlib.h>

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.75;
void __global__ add(const double* x, const double* y, double* z,int N);
void check(const double* z, const int N);

int main(void) {
	//Host 变量，数组并进行初始化
	const int N = 1000;
	const int M = sizeof(double) * N;
	double* h_x = (double*)malloc(M);
	double* h_y = (double*)malloc(M);
	double* h_z = (double*)malloc(M);

	for (int n = 0; n < N; ++n) {
		h_x[n] = a;
		h_y[n] = b;
	}

	//Device，定义3个数组并分配了内存(显存)  malloc动态分配内存
	// 用cudaMalloc(变量名，大小)函数分配到GPU中。
	//函数原型 cudaError_t cudaMalloc(void **address, size_t size);
	//address时待分配设备内存的指针，因为内存本省就是一个指针，所以为指针的指针，双重指针。
	//size待分配内存的字节数
	//返回值是一个错误代号。如果调用成功，返回cudaSuccess，否则返回错误。
	//该函数为某个变量分配size字节的线性内存（linearmemory）
	//d_x 是一个double 类型的指针，那么它的地址&d_x就是double类型的双重指针。
	//而(void **) 是一个强制类型转换操作，将一个某种类型的双重指针转换为一个void类型的双重指针。
	//这种类型转换可以不明确地写出来，简写为，cudaMalloc(&d_x, M);
	//该函数的功能是改变指针d_x本身的值（将一个指针赋值给d_x），而不是改变d_x所指内存缓冲区中的变量值。
	//函数cudaMalloc()要求用传双重指针的方式改变一个指针的值，而不是直接返回一个指针，是因为该函数已经将返回值用于返回错误代号，而C++又不支持多个返回值。
	//cudaMalloc()函数可以为不同类型的指针变量分配设备内存。
	double* d_x, * d_y, * d_z;
	cudaMalloc((void**)&d_x, M);
	cudaMalloc((void**)&d_y, M);
	cudaMalloc((void**)&d_z, M);
	//为了区分主机变量和设备变量，d作为设备变量，h作为主机变量。
	//cudaError_t cudaMemcpy (void  *dst, const void *src, size_t count,enum cudaMemcpyKind kind) 
	//dsr 目标地址，src源地址，conut，复制数据的字节数
	//kind，枚举类型变量，标志数据传递方向
	/*
	cudaMemcpyHostToHost，表示从主机复制到主机。
	– cudaMemcpyHostToDevice，表示从主机复制到设备。
	– cudaMemcpyDeviceToHost，表示从设备复制到主机。
	– cudaMemcpyDeviceToDevice，表示从设备复制到设备。
	– cudaMemcpyDefault，表示根据指针dst和src所指地址自动判断数据传输的方向。
	 该函数的作用是将一定字节数的数据从源地址所指缓冲区复制到目标地址所指缓冲区
	*/

	cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
	cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);
	
	//核函数执行配置参数
	const int block_size = 128; //128个一维线程块
	const int grid_size = (N+block_size-1) / block_size; //10^8/128个线程块
	//调用核函数计算
	add << <grid_size, block_size >> > (d_x, d_y, d_z,N);

	cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
	check(h_z, N);

	//当然在主机中也可以用new，和delete使用。
	free(h_x);
	free(h_y);
	free(h_z);

	//cudaMalloc也需要cudaFree来释放内存变量。
    //cudaError_t cudaFree(void* address);
	//参数address就是待释放的设备内存变量（不是双重指针）
	//变量释放要一一对应，不能把设备变量用c++free给释放了
	cudaFree(d_x);
	cudaFree(d_y);
	cudaFree(d_z);

	return 0;

}

//核函数要求，返回类型必须时void，可以用return关键字，但不可返回任何值
//必须使用限定词 __global__,也可以加static，限定符的顺序可以任意
//函数名无特殊要求，冰球支持重载，即可以用同一颗函数名，表示具有不同参数列表的函数。
//不支持可变数量的参数列表，即参数的个数必须确定
//可以向核函数传递非指针变量，其内容对每个线程可见。
//除非使用统一内存编程机制，否则传给核函数的数组（指针）必须只想设备内存。
//核函数不可成为一个类的成员。通常做法时用一个包装函数调用核函数，而将包装函数定义为类的成员。
//从计算能力3.5开始，引入了动态并行机制，在核函数内部可以调用其他核函数，甚至可以调用自己（即递归）。
//无论是从主机调用还是从设备调用，核函数都是在设备中执行。调用核函数时必须指定执行配置，即三括号和它里面的参数。


void __global__ add(const double* x, const double* y, double* z,int N) {
	//核函数中数据与线程的对应
	//单指令-多线程
	//n 表示数组元素指标
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if (n < N) {
		z[n] = x[n] + y[n];
	}

}

void check(const double* z, const int N) {
	bool has_error = false;
	for (int n = 0; n < N; ++n) {
		if (fabs(z[n] - c) > EPSILON) {
			has_error = true;
		}
	}
	printf("%s\n", has_error ? "Has error" : "No errors");

}

上述代码运行结果：

这个我也不太清楚，应该输出has error，感觉应该是我电脑GPU精度不高。

此外也尝试了用c写的，还是精度有问题，不太明白，有懂得大佬可以帮忙解释下。

这里我问了同学，double双精度的误差存在是正常的。主要还是看这个代码能否正常运行。

双精度一般是63位，包括符号位，11位指数位和52个尾数位。精度和范围更高。

一些注释和必要的内容我已经加在代码里了，可以学习。

总结一下吧，

cudaMalloc（），为变量分配设备内存，需要配合cudaFree（）来将内存释放

cudaMemcpy，可以进行主机与设备之间的数据传递，此外主机到主机，设备到主机，设备到设备也都行。

核函数中需要计算数据的下标，将数据对应，然后进行运算，数据的下标计算在上一节中有提到。

核函数有一定要求，限定词__global__不可少，返回值类型必须是void

核函数的配置必须有,网格大小，线程块大小。

核函数中的if语句必要性：

代码中的核函数没有使用参数N，当N是blockDim.x即（block_size）的整数倍时，不会引起问题，因为核函数中的线程数目刚好等于数组元素的个数。然而当不是整数倍时就有可能发生错误。

来自cuda编程与实践

总结即，就是当多余一些数据时，需要加128，线程，然后用if条件语句让没有数据的线程不运行计算过程。代码中已经修改过了。

3.自定义设备函数

设备函数：

核函数可以调用不带执行配置的自定义函数，这样的自定函数称为设备函数（device function）。他是在设备中执行，并在设备中被调用的。

与之相比，核函数时在设备中执行，但在主机端被调用。现在也支持在一个核函数中调用其他和函数，甚至该核函数本身。

函数执行空间表示符：

__gloabal__:核函数，一般由主机调用，在设备中执行，如果使用动态并行，则也可以在核函数中调用自己或其他核函数。

__device__:设备函数，只能被核函数或其他设备函数调用，在设备中执行。

__host__:主机端普通C++函数，在主机中被调用，在主机中执行。对于主机端函数，该修饰符可以省略。之所以提供这样一个修饰符，是因为有时可以用host和device同时修饰一个函数，是的该函数既是一个c++中的普通函数，又是一个设备函数。这样做可以减少冗余代码。编译器将正对主机和设备分别编译改函数。

不能同时使用device和global修饰一个函数，即不能将一个函数同时定义为设备函数和核函数。

也不能同时用host和global修饰一个函数，即不能将一个函数同时定义为主机函数和核函数。

编译器决定把设备函数当作内联函数（inline function）或非内联函数，但可以用修饰符__noinline__建议一个设备函数为非内联函数，也可以用修饰符__forceinline__建议一个设备函数为内联函数。

什么是内联函数：

内联函数是一种特殊类型的函数，它的目的是提高程序的执行效率。内联函数在编译时会被直接嵌入到调用它的地方，而不是像普通函数一样在运行时进行函数调用和返回。这减少了函数调用的开销，因为没有函数栈帧的创建和销毁，但也增加了代码的大小。内联函数通常用于小而频繁调用的函数，以减少函数调用的开销。

在使用内联函数时，需要注意以下几点：

1. 内联函数通常适用于简单的函数，不适用于复杂的函数。
2. 内联函数的定义通常应该放在头文件中，以便多个源文件都能访问它。
3. 编译器有权忽略内联函数的内联请求，因此inline关键字只是一个建议，最终是否内联由编译器决定。
4. 内联函数不应该包含递归调用或包含复杂的控制流，因为这会导致生成的代码变得复杂，甚至比函数调用更慢。
5. 内联函数的过度使用可能会导致代码膨胀，因此需要权衡内联和代码大小之间的关系。

总结，内联函数增加代码大小，但在编译时直接嵌入到调用它的地方，可以减少函数调用的开销。内联函数适合简单函数，inline只是一个建议，编译器回决定是否内联。不能过度使用内联函数，会使得代码量膨胀。

例子，为数组相加的核函数定义一个设备函数：

3个版本，返回值，指针，引用：

//有返回值的设备函数
double __device__ add1_device(const double x, const double y) {
	return(x + y);
}

void __global__ add1(const double* x, const double* y, double* z,const int N) {
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if (n < N) {
		z[n] = add1_device(x[n],y[n]);
	}
}

//用指针的设备函数
double __device__ add2_device(const double x, const double y,double *z) {
	*z = x + y;
}

void __global__ add2(const double* x, const double* y, double* z, const int N) {
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if (n < N) {
		 add2_device(x[n], y[n],&z[n]);
	}
}

//用引用的设备函数
double __device__ add3_device(const double x, const double y, double& z) {
	z = x + y;
}

void __global__ add3(const double* x, const double* y, double* z, const int N) {
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	if (n < N) {
		add3_device(x[n], y[n], z[n]);
	}
}

设备函数被核函数调用。