03《CUDA编程基础与实践》-简单CUDA程序的基本框架

第3章 简单CUDA程序的基本框架

3.1 例子：数组相加

#include<math.h>
#include<stdlib.h>
#include<stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
void add(const double* x, const double* y, double* z, const int N);
void check(const double* z, const int N);

int main()
{
	const int N = 100000000;
	const int M = sizeof(double) * N;
	double* x = (double*)malloc(M);
	double* y = (double*)malloc(M);
	double* z = (double*)malloc(M);
	for (int n = 0; n < N; n++) {
		x[n] = a;
		y[n] = b;
	}
	add(x, y, z, N);
	check(z, N);
	free(x);
	free(y);
	free(z);
	return 0;
}

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 errors" : "No errors");
}
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];
	}
}

3.2 CUDA 程序的基本框架

#include<stdio.h>
#include<cuda.h>
#include <cuda_runtime_api.h>
#include<device_launch_parameters.h>
#include<math.h>
#include <cassert>
const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
void __global__ add(const double* x, const double* y, double* z);
void check(const double* z, const int N);
int main() {
	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);
	for (int i = 0; i < N; i++) {
		h_x[i] = a;
		h_y[i] = b;
	}
	double* d_x, * d_y, * d_z;
	cudaMalloc((void**)&d_x, M);
	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;
	add<<<grid_size, block_size>>>(d_x, d_y, d_z);
	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(const double* x, const double* y, double* z) {
	const int n = blockDim.x * blockIdx.x + threadIdx.x;
	z[n] = x[n] + y[n];
}
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 ? "有错" : "无错");
}

3.3自定义设备函数

核函数可以调用不带执行配置的自定义函数，这样的自定义函数称为设备函数(device function)。它是在设备中执行，并在设备中被调用的。与之相比，核函数是在设备中执行，但在主机端被调用的。现在也支持在一个核函数中调用其他核函数，甚至调用该核函数本身，但本书不涉及这方面的内容。设备函数的定义与使用涉及 CUDA 中函数执行空间标识符的概念。我们先对此进行介绍，然后以数组相加的程序为例展示设备函数的定义与调用。

3.3.1 函数执行空间标识符

在CUDA 程序中，由以下标识符确定一个函数在哪里被调用，以及在哪里执行:
(1)用__global__修饰的函数称为核函数，一般由主机调用，在设备中执行如果使用动态并行，则也可以在核函数中调用自己或其他核函数。
(2)用 device__修饰的函数称为设备函数，只能被核函数或其他设备函数调用，在设备中执行。
(3)用__host__修饰的函数就是主机端的普通 C++函数，在主机中被调用在主机中执行。对于主机端的函数，该修饰符可省略。之所以提供这样一个修饰符是因为有时可以用__host__和__device 同时修饰一个函数，使得该函数既是一个 C++ 中的普通函数，又是一个设备函数。这样做可以减少冗余代码。编译器将针对主机和设备分别编译该函数。
(4)不能同时用 __device__和__global__修饰一个函数，即不能将一个函数同时定义为设备函数和核函数。
(5)也不能同时用 __host__和 __global__修饰一个函数，即不能将一个函数同时定义为主机函数和核函数。
(6)编译器决定把设备函数当作内联函数(inline function)或非内联函数，但可以用修饰符__noinline__建议一个设备函数为非内联函数(编译器不一定接受)也可以用修饰符 __forceinline__建议一个设备函数为内联函数。