CUDA编程极简入门（二）：Hello World

近年来，随着人工智能的发展，“卖铲人”NVIDIA在人工智能行业的地位越来越重要。CUDA编程也是人工智能重要知识点，本系列通过学习三章文章对CUDA编程有个快速理解和入门：第一章介绍基础知识点，理解CUDA编程基础概念；第二章通过2个简单程序的编程、编译和执行对CUDA编程有个初步掌握；第三章通过CUDA并行编程的实操掌握入门。

CUDA 与 C 的快速比较

下表并排比较了 C 和 CUDA 中的 hello world 程序。

C

void c_hello(){
    printf("Hello World!\n");
}

int main() {
    c_hello();
    return 0;
}

CUDA

__global__ void cuda_hello(){
    printf("Hello World from GPU!\n");
}

int main() {
    cuda_hello<<<1,1>>>(); 
    return 0;
}

C 和 CUDA 实现的主要区别在于global说明符和<<<…>>>语法。global说明符表示在设备（GPU）上运行的函数。此类函数可以通过主机代码调用，例如main()示例中的函数，也称为“内核”。

调用内核时，其执行配置通过<<<…>>>语法提供，例如cuda_hello<<<1,1>>>()。在 CUDA 术语中，这称为“内核启动(1,1)”。我们将在第三章讨论该参数。

编译 CUDA 程序

编译 CUDA 程序与编译 C 程序类似。NVIDIA nvcc 在 CUDA 工具包中提供了一个名为 CUDA 编译器来编译 CUDA 代码，通常存储在扩展名为 的文件中 .cu 。例如

$> nvcc hello.cu -o hello

使用上述命令编译 CUDA 程序时，你可能会看到以下警告

nvcc warning : The 'compute_20', 'sm_20', and 'sm_21' architectures are deprecated, and may be removed in a future release (Use -Wno-deprecated-gpu-targets to suppress warning).

现在可以忽略此警告。

执行

CUDA hello world 示例不执行任何操作，即使程序已编译，屏幕上也不会显示任何内容。为了使操作生效，我们将研究向量加法。

以下是用 C 语言实现的向量加法示例（ ）。该示例计算存储在数组和./vector_add.c中的两个向量的加法，并将结果放入数组 中。a b out

#define N 10000000

void vector_add(float *out, float *a, float *b, int n) {
    for(int i = 0; i < n; i++){
        out[i] = a[i] + b[i];
    }
}

int main(){
    float *a, *b, *out; 

    // Allocate memory
    a   = (float*)malloc(sizeof(float) * N);
    b   = (float*)malloc(sizeof(float) * N);
    out = (float*)malloc(sizeof(float) * N);

    // Initialize array
    for(int i = 0; i < N; i++){
        a[i] = 1.0f; b[i] = 2.0f;
    }

    // Main function
    vector_add(out, a, b, N);
}

练习：将向量加法转换为 CUDA

在第一个练习中，我们将使用 hello world 作为示例转换 vector_add.c 为 CUDA 程序。 vector_add.cu

复制vector_add.c到****vector_add.cu

$> cp vector_add.c vector_add.cu

转换vector_add()为 GPU 内核

__global__ void vector_add(float *out, float *a, float *b, int n) {
    for(int i = 0; i < n; i++){
        out[i] = a[i] + b[i];
    }
}

vector_add()将调用更改main()****为内核调用

vector_add<<<1,1>>>(out, a, b, N);

编译并运行程序

$> nvcc vector_add.c -o vector_add
$> ./vector_add

你会注意到程序无法正常工作。原因是 CPU 和 GPU 是独立的实体。两者都有自己的内存空间。CPU 不能直接访问 GPU 内存，反之亦然。在 CUDA 术语中，CPU 内存称为主机内存，GPU 内存称为设备内存。指向 CPU 和 GPU 内存的指针分别称为主机指针和设备指针。

设备内存管理

CUDA 提供了几个用于分配设备内存的函数。最常见的是 cudaMalloc() 和 cudaFree() 。这两个函数的语法如下

cudaMalloc(void **devPtr, size_t count);
cudaFree(void *devPtr);

cudaMalloc() count在设备内存中分配大小的内存，并将设备指针更新devPtr为分配的内存。cudaFree()释放设备指针指向的设备内存区域<font style="color:rgb(231, 76, 60);background-color:rgb(252, 252, 252);">devPtr</font>。它们分别相当于 C 中的malloc()和free()

内存传输

主机和设备内存之间的数据传输可以通过函数完成 cudaMemcpy ，类似于 memcpy C 语言。语法 cudaMemcpy 如下

cudaMemcpy(void *dst, void *src, size_t count, cudaMemcpyKind kind)

函数将大小为 的内存count从src复制到dst。kind表示方向。，对于典型用法， kind的值是cudaMemcpyHostToDevice或cudaMemcpyDeviceToHost。 还有其他可能的值，但在本教程中我们不会涉及它们。

练习（续）：完成向量加法

1. 为数组 a 、 b 和分配和释放设备内存 out 。 2. 在主机和设备内存之间传输 a 、 b 和 out 。 - 测验：内核执行之前和之后必须传输哪个数组？

示例：数组“a”的解决方案

```plain void main(){ float *a, *b, *out; float *d_a;

a = (float*)malloc(sizeof(float) * N);

// Allocate device memory for a
cudaMalloc((void**)&d_a, sizeof(float) * N);

// Transfer data from host to device memory
cudaMemcpy(d_a, a, sizeof(float) * N, cudaMemcpyHostToDevice);

…
vector_add<<<1,1>>>(out, d_a, b, N);
…

// Cleanup after kernel execution
cudaFree(d_a);
free(a);

}

1. <font style="color:rgb(64, 64, 64);background-color:rgb(252, 252, 252);">编译并测量性能。（请参阅（</font>[<font style="color:rgb(64, 64, 64);background-color:rgb(252, 252, 252);">./solutions/vector_add.cu</font>](https://cuda-tutorial.readthedocs.io/en/latest/tutorials/tutorial01/solutions/vector_add.cu)<font style="color:rgb(64, 64, 64);background-color:rgb(252, 252, 252);">）中的解决方案）</font>

```plain
$> nvcc vector_add.cu -o vector_add
$> time ./vector_add

性能分析

使用 time 不会提供有关程序性能的太多信息。NVIDIA 提供了一个名为 的命令行分析器工具 nvprof ，它提供了有关 CUDA 程序性能的更多深入信息。

要分析向量加法，请使用以下命令

$> nvprof ./vector_add

以下是 Tesla M2050 上的分析结果示例

==6326== Profiling application: ./vector_add
==6326== Profiling result:
Time(%)      Time     Calls       Avg       Min       Max  Name
 97.55%  1.42529s         1  1.42529s  1.42529s  1.42529s  vector_add(float*, float*, float*, int)
  1.39%  20.318ms         2  10.159ms  10.126ms  10.192ms  [CUDA memcpy HtoD]
  1.06%  15.549ms         1  15.549ms  15.549ms  15.549ms  [CUDA memcpy DtoH]

总结

在本章中，我们演示了如何在 CUDA 中编写一个简单的向量加法。我们介绍了 GPU 内核及其从主机代码的执行。此外，我们还介绍了 CPU 和 GPU 之间分离内存空间的概念。我们还演示了如何管理设备内存。

但是，我们仍然没有并行运行程序。内核执行配置<<<1,1>>>表明内核仅以 1 个线程启动。在下一章中，我们将修改向量加法以并行运行。