说在前面：本文主要给出一个简单的CUDA程序，用来说明一个完整的核函数是如何在CPU端进行调用的，并且会对相关代码做出详细解释。
一. 完整代码
在这篇文章中我们给出了一个用CUDA核函数实现向量加法的代码，并且完成了CPU端对核函数的调用。
#include <stdio.h>
#include <cuda.h>
#include <cuda_runtime.h>
typedef float FLOAT;
/* CUDA kernel function */
__global__ void vec_add(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
/* 1D grid */
int idx = blockDim.x * blockIdx.x + threadIdx.x;
if (idx < N) z[idx] = y[idx] + x[idx];
}
void vec_add_cpu(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
for (int i = 0; i < N; i++) z[i] = y[i] + x[i];
}
int main()
{
int N = 10000;
int nbytes = N * sizeof(FLOAT);
/* 1D block */
int bs = 256;
/* 1D grid */
int s = ceil((N + bs - 1.) / bs);
dim3 grid(s);
FLOAT *dx, *hx;
FLOAT *dy, *hy;
FLOAT *dz, *hz;
/* allocate GPU mem */
cudaMalloc((void **)&dx, nbytes);
cudaMalloc((void **)&dy, nbytes);
cudaMalloc((void **)&dz, nbytes);
/* init time */
float milliseconds = 0;
/* alllocate CPU mem */
hx = (FLOAT *) malloc(nbytes);
hy = (FLOAT *) malloc(nbytes);
hz = (FLOAT *) malloc(nbytes);
/* init */
for (int i = 0; i < N; i++) {
hx[i] = 1;
hy[i] = 1;
}
/* copy data to GPU */
cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
/* launch GPU kernel */
vec_add<<<grid, bs>>>(dx, dy, dz, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&milliseconds, start, stop);
/* copy GPU result to CPU */
cudaMemcpy(hz, dz, nbytes, cudaMemcpyDeviceToHost);
/* CPU compute */
FLOAT* hz_cpu_res = (FLOAT *) malloc(nbytes);
vec_add_cpu(hx, hy, hz_cpu_res, N);
/* check GPU result with CPU*/
for (int i = 0; i < N; ++i) {
if (fabs(hz_cpu_res[i] - hz[i]) > 1e-6) {
printf("Result verification failed at element index %d!\n", i);
}
}
printf("Result right\n");
printf("Mem BW= %f (GB/sec)\n", (float)N*4/milliseconds/1e6);///1.78gb/s
cudaFree(dx);
cudaFree(dy);
cudaFree(dz);
free(hx);
free(hy);
free(hz);
free(hz_cpu_res);
return 0;
}二. 代码详解
2.1. 头文件
这个代码中包含三个头文件stdio.h，cuda.h，cuda_runtime.h
-
这是一个标准 C 语言库的头文件，用于提供输入和输出功能，如打印到控制台。它不是专门为 CUDA 设计的，而是 C 语言的一部分。在 CUDA 程序中使用这个头文件通常是为了在主机端（CPU 端）执行如#include <stdio.h>:printf等常规 I/O 操作。 -
这个头文件是 CUDA 驱动 API 的一部分。CUDA 驱动 API 提供了低级别的 CUDA 功能。它允许更细粒度的控制，包括设备初始化、上下文管理、内存分配等。由于其低级性质，这个 API 对于需要精细控制 CUDA 资源的高级用户更加适用。#include <cuda.h>: -
#include <cuda_runtime.h>:cuda_runtime.h是 CUDA 运行时 API 的头文件。CUDA 运行时 API 是一个更高级别的抽象，提供了更易于使用的接口，用于管理设备内存、执行核函数等。对于大多数 CUDA 应用程序来说，这个 API 提供了足够的功能和简化的接口，使得编程更加容易。它处理很多底层细节，如自动的设备上下文管理。
2.2. vec_add函数
__global__ void vec_add(FLOAT *x, FLOAT *y, FLOAT *z, int N)这是一个cuda核函数，__global__用于标记这是一个再gpu上执行的核函数，这里的x和y代表输入的需要执行向量加法的数据，z用于存储加法的结果，N表示向量的长度是N。
int idx = blockDim.x * blockIdx.x + threadIdx.x;关于这行代码，我们在之前的博客中有提到过，blockDim.x表示每个block在x方向上的线程数量，blockIdx表示当前block在grid的x方向上的相对id，threadIdx.x表示当前线程在当前block的x方向上的相对id。可以看出idx就是当前线程在线程分布中的全局id。另外需要注意，可以看出这里的grid是一维的。
if (idx < N) z[idx] = y[idx] + x[idx];这里idx小于N，是因为我们需要处理的向量长度就是N，那么idx >=N的线程就不需要执行任何操作。z[idx] = y[idx] + x[idx]就是对应的向量加法，每个线程执行向量中一个元素的加法操作。
2.3. vec_add_cpu函数
void vec_add_cpu(FLOAT *x, FLOAT *y, FLOAT *z, int N)
{
for (int i = 0; i < N; i++) z[i] = y[i] + x[i];
}这个函数就是普通的cpu上执行的函数，但是它是单线程执行整个向量的加法，所以在执行效率上是不如vec_add函数的。
2.4. main函数
int N = 10000;
int nbytes = N * sizeof(FLOAT);
/* 1D block */
int bs = 256;
/* 1D grid */
int s = ceil((N + bs - 1.) / bs);
dim3 grid(s);
FLOAT *dx, *hx;
FLOAT *dy, *hy;
FLOAT *dz, *hz;
/* allocate GPU mem */
cudaMalloc((void **)&dx, nbytes);
cudaMalloc((void **)&dy, nbytes);
cudaMalloc((void **)&dz, nbytes);
/* init time */
float milliseconds = 0;
/* alllocate CPU mem */
hx = (FLOAT *) malloc(nbytes);
hy = (FLOAT *) malloc(nbytes);
hz = (FLOAT *) malloc(nbytes);
/* init */
for (int i = 0; i < N; i++) {
hx[i] = 1;
hy[i] = 1;
}这里的N=10000表示我们的两个向量长度是10000，bs=256表示每个blcok的线程数是256个，s是执行 CUDA 内核所需的最小网格大小，确保有足够的线程来处理所有的数据元素。cudamalloc用于分配device端内存，包括两个输入dx，dy和一个输出dz，malloc用于分配host端内存，包括hx，hy和hz，hx和hy分别指向两个需要计算的向量。最后的for循环是对host端输入数据的简单初始化。
cudaMemcpy(dx, hx, nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, nbytes, cudaMemcpyHostToDevice);cudaMemcpy函数是一个cuda运行时api，用于数据在cpu和gpu上的迁移，上面是将两个向量从host端转移到device端，需要注意的是数据的移动也是影响程序执行效率的一个重要方面。
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start);
/* launch GPU kernel */
vec_add<<<grid, bs>>>(dx, dy, dz, N);
cudaEventRecord(stop);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&milliseconds, start, stop); 这段代码使用 CUDA 事件来测量 GPU 内核函数的执行时间。CUDA 事件提供了一种精确测量在 GPU 上发生操作时间的方式。下面是这段代码的详细解释：
-
创建 CUDA 事件：
cudaEvent_t start, stop;: 定义两个 CUDA 事件变量start和stop。cudaEventCreate(&start);: 创建start事件。cudaEventCreate(&stop);: 创建stop事件。
这些事件将用于标记内核执行的开始和结束时间。
-
记录开始时间：
cudaEventRecord(start);: 记录start事件。这标志着后续操作（如内核执行）的开始时间。 -
启动 GPU 内核：
vec_add<<<grid, bs>>>(dx, dy, dz, N);: 这是一个 CUDA 内核函数调用。vec_add是内核函数名，grid和bs分别是网格和块的大小，dx,dy,dz,N是传递给内核的参数。 -
记录结束时间：
cudaEventRecord(stop);: 立即在内核调用之后记录stop事件。这标志着内核操作结束的时间。 -
同步事件：
cudaEventSynchronize(stop);: 这个调用确保stop事件已经发生，即内核操作已经完成。这是必要的，因为 CUDA 内核调用是异步的，cudaEventRecord不会等待内核完成。 -
计算经过的时间：
cudaEventElapsedTime(&milliseconds, start, stop);: 计算从start到stop事件之间的时间，以毫秒为单位。结果存储在milliseconds变量中。
这种方法是测量 CUDA 内核执行时间的一种标准方式。它提供了一种相对精确的方法来衡量 GPU 上函数执行的时间。
cudaMemcpy(hz, dz, nbytes, cudaMemcpyDeviceToHost);在执行完核函数后，将gpu端获得的结果转移到cpu端。
/* CPU compute */
FLOAT* hz_cpu_res = (FLOAT *) malloc(nbytes);
vec_add_cpu(hx, hy, hz_cpu_res, N);
这段代码是通过调用vec_add_cpu调用cpu计算向量hx和hy相加，结果保存在hz_cpu_res中。
/* check GPU result with CPU*/
for (int i = 0; i < N; ++i) {
if (fabs(hz_cpu_res[i] - hz[i]) > 1e-6) {
printf("Result verification failed at element index %d!\n", i);
}
}
printf("Result right\n");
printf("Mem BW= %f (GB/sec)\n", (float)N*4/milliseconds/1e6);///1.78gb/s比较cpu和gpu计算的向量相加的结果是否一致。
cudaFree(dx);
cudaFree(dy);
cudaFree(dz);
free(hx);
free(hy);
free(hz);
free(hz_cpu_res);释放资源，cpu资源释放使用free，gpu资源释放使用cudafree。
三. 总结一下
上面我们给出了一个cpu--gpu--cpu的完整代码，并且给出了比较详细的解释。一个完整的程序应该包含下面几个方面：
1. host端申请内存并初始化数据，host端也就是cpu端
2. device端申请内存，device端也就是gpu端，申请的内存也就是显存
3. 将host端的数据拷贝到device端
4. 执行CUDA核函数，在下面这个例子中的核函数就是vec_add
5. 将数据在gpu上处理完后，再将device端的结果拷贝回host端
6. 通常还需要将host端计算结果和device端的计算结果进行比较，验证结果误差
7. 释放之前host端和device端的申请的内存
小知识
c++中的typedef和#define的区别：
typedef 是 C/C++ 语言的一部分，用于为类型创建新名称。它通常用于简化复杂类型的声明或提高代码的可读性。例如，你可以使用 typedef 为结构体或指针类型定义一个新名字。typedef 仅限于类型定义。
#define 是预处理指令，用于定义宏。宏可以是一个值、一个代码片段或者几乎任何你希望在代码中多次出现的东西。预处理器在编译之前将所有的宏替换为它们的定义。
502
被折叠的 条评论
为什么被折叠?
到【灌水乐园】发言
