CUDA入门必看，如何高效地编写并行程序

CUDA入门必看，如何高效地编写并行程序

本文链接

进入公司实习已经一个月有余，从编写第一个 kernel 开始到现在分析优化 LLM 程序，我的CUDA学习成果颇丰，项目进展顺利，现将我的学习路径整理分享出来。跟随在GPU芯片架构领域深耕多年的企业家王旭老师，我从一开始对GPU架构领域全无了解，到如今上手LLM开源项目优化项目，在该领域我逐渐从一个的小白成长为资深小白。从一个企业级程序架构师的角度，王旭老师为我整理了一套CUDA学习方案和文档资料。

虽然CUDA架构是基于C++语言所开发的，但它不仅仅只是在GPU调用、线程分配方向上API的延伸拓展，而是遵循着一定的设计思路。CUDA的研发以及在当下的流行，从始至终都在完成提升性能这一件事。从这一套学习方案中，你会时刻感受到性能指标在编写CUDA程序中所占据的考量有多重。因此要编写出更加成熟的kernel程序，可以跟着我梳理出的学习流程一步一步掌握基础知识，侧重于对性能指标的提升，将CUDA最开始的初衷运用到你的代码之中。

要学习它最前沿的知识，就先了解它最初的样貌。

学习CUDA我推荐先从Nvidia CUDA 官方文档开始。我将我的学习路线顺序整理如下，并对每一章节的知识作出归纳和总结，并附上我对初学者所给出的建议，便于你理解每一篇文档的用意，所要达到的目的，帮助你快速高效地攻克在学习CUDA道路上的重重难点，获得知识的同时在理念上的有所感悟。

我将不对代码细节上做过多的讲述，如果你需要的是保姆级教程, 也许这篇文档并不适合你。但是如果你在细节上遇到许多坑，那么这篇文档将为你提供解决问题的思路。

1.Cuda 简介

主题：分配GPU线程 & 调度GPU内存

原文链接: An Even Easier Introduction to CUDA

如果你在耕地，你需要什么？一头壮牛还是1024只鸡？

你只需要在 main 函数中敲入这一行简单的代码，就能体会到调用GPU函数是一件多么自然的事。

kernel<<<grid, block>>>(char* param);

其中kernel是你的函数名，你可以自定义。在初学阶段，你可以将grid设置为 1 ，代表你只用到一个block线程块。block是你的线程块尺寸，代表着一个线程块中线程thread的多少。grid和block可以是一维的，也可以是二维的，甚至是三维的，取决于你要完成的任务。在上图的例子中，你可以将block简单设置为256。你知道它最好是2的幂次(但<=1024)，它不仅仅只是程序员的强迫症。

在你的kernel函数实现前加上 __global__,代表它是运行在GPU的函数。并且你的每一个线程中都在调用这个函数。

__global__ kernel(char* param) { ... }

为了能够在不同的线程中处理不同的数据，假设现在有一块内存地址，我需要在线程中拿到这些内存地址，访问它的数据。那么我们这样做，给每个线程分配一个id序号，它们从0 ~ n，n是id数最大的线程，其实在一维线程块中它等于上面的block - 1，即255。我们用这些序号的差异来计算得出每个线程需要访问的地址。

在CUDA语言中，我们一般称GPU为设备Device，称CPU为主机Host。假设我们的Device上有一块内存 d_mem , 我们这样来访问到它的第 0 ~ 255 个数据。

__global__ kernel(int* d_mem) { 

	int id =  treadIdx.x;
	d_mem[id] = 0;

}

这个函数在每一个线程中都会被调用。此时想必你已经知道了如何使用线程的序号来控制你想要访问的内存位置。下面的几个API是我简单引用的最基础的API。

blockIdx.x  //此线程的block序号
blockDim.x  //block在x维上的长度，在一维block中代表着一个线程块中的thread数目
gridDim.x   //grid在x维上的长度，在一维grid中代表着一个grid中的block数目

经过学习，你会从一维的block/grid逐步接触到二维、三维的block/grid，但现在还不必要。请认真思考这些变量的含义，并把它们抽象成图格的形式，这样你会更加理解其中的含义。

上图中的index的含义是什么？以后你将经常用到这样的形式去处理更多的数据。index = blockIdx.x * blockDim.x + threadIdx.x。这里我需要做一点提醒，上图中的gridDim.x是线程块block的长度，也就是说有 4096 个 0~255 那么长的block。曾经我误以为这是 thread的数目。

---

2.GPU-CPU共用内存：统一内存

主题：统一内存调度 & 内存创建销毁、拷贝

原文链接: Unified Memory for CUDA Beginners

程序员在某种意义上也属于搬运工。

当你需要使用线程处理数据的时候，一定离不开内存空间的创建。在普通的C语言程序中，我们只需要调用动态分配内存的函数 malloc 就能在内存中开辟一块指定大小的内存空间。但是 malloc 所创建的内存空间是 Host 端的，你的 Decive 一般来说是无法访问的。 CUDA 中提供了一个创建统一内存的API, 这个API的使用和 malloc 类似，但是创建的内存可以被 device 端调用。也就是说，你的__global__函数可以访问这块内存。

float* x;
float* y;

cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));

作为初学者你可以先使用它，而不去思考这块内存究竟被放置在哪个位置。 你可以通过文档继续了解它的实现原理，简单来说，尽管这个函数在 Host 和 device 上都创建了内存，但是只有当你需要在 device 上使用时，它才会自动将 Host 上的内存拷贝到 device 上。

以下是三个创建内存的API。

malloc()            //原生C语言的动态内存创建函数
cudamalloc()        //特别的，只在 Device 端创建内存
cudamallocHost()    //特别的，只在 Host 端创建内存

你可以使用 malloc() 或者 cudamallocHost() 在 Host 端创建内存，他们创建的内存在传输的过程中有所不同。有关这种不同，我会在后面的“数据传输”小节简单为你介绍。在现阶段的学习中，你可以随意使用二者其一。在细节上，cudamallocHost() 的使用语法与 cudamalloc() 相似，从以下的调用示例中你就能发现。 以下的示例中展示了在 Host 端创建大小为 bytes 字节的内存，储存的数据类型为 float 。

float *h_malloc, *h_cudamallocHost;    //声明指针型变量

h_malloc = (float *)malloc(bytes);                    //使用 malloc 创建 Host 内存
cudaMallocHost((void **)&h_cudamallocHost, bytes);    //使用 cudamallocHost 创建 Host 内存

创建好内存，你需要对你的数据进行初始化， memset() 函数很有帮助。 想起曾经大一课堂上老师为我们介绍这类函数的用法，他以过来人告诉我们这类函数在实际编程经常使用到。 作为大一的愣头青，心是向着算法的，自然没把搬运内存和拷贝内存的事放在心上。死去的记忆开始攻击我(ー`´ー)，如今每天都在开辟内存、初始化内存、拷贝内存、再把内存拷贝回来。以下是几个内存拷贝的函数，memcpy()想必你已经不陌生, cudamemcpy()函数是cuda的内存拷贝API，其中第四个参数决定了内存从哪拷贝到哪。

memcpy(*dest, *src, byteSize)  //从src地址拷贝byteSize字节大小的内存给dest的地址
cudamemcpy(*dest, *src, byteSize, cudamemcpyHostToDevice); //从Host拷贝给Device
cudamemcpy(*dest, *src, byteSize, cudamemcpyDeviceToHost); //从Device拷贝回Host

如果你使用完刚才开辟的内存后顺手把它释放了，那么你是一个好程序员。

Free(*mem);      //释放由 malloc 创建的内存
cudaFree(*mem);  //释放由 cudamalloc() cudamallocHost() 创建的内存

你已经是一个成熟的程序员了，现在你要完成组织交给你的任务ヾ(≧▽≦*)o。以下是一段完整的CUDA代码，其中的saxpy是运行在GPU上的kernel函数，它可以将 x, y 矩阵的对应的元素做一次乘加操作。 将这段代码复制到你的.cu文件里，通过编译运行它来直接感受kernel函数是如何被调用的，并了解我们是怎么在 Host 和 device 端进行数据传输的，我想这会对你有帮助。

#include <stdio.h>

__global__ void saxpy(int n, float a, float *x, float *y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        y[i] = a * x[i] + y[i];
}

int main(void)
{
    int N = 20 * (1 << 20);
    float *x, *y, *d_x, *d_y;
    x = (float *)malloc(N * sizeof(float));
    y = (float *)malloc(N * sizeof(float));

    cudaMalloc(&d_x, N * sizeof(float));
    cudaMalloc(&d_y, N * sizeof(float));

    for (int i = 0; i < N; i++)
    {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);

    saxpy<<<(N + 511) / 512, 512>>>(N, 2.0f, d_x, d_y);

    cudaMemcpy(y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost);

    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
    {
        maxError = max(maxError, abs(y[i] - 4.0f));
    }
    printf("Max error: %fn\n", maxError);

}

这段代码在主程序中添加了一个循环用于判断计算是否全部正确，如果计算全部正确，会打印"Max error: 0.0"的输出结果。

其中的 (N + 511) / 512 运算是否让你费解？实际上，根据我们想要处理的数据量，例如上述代码的 20M 数据，我们给每个 float 数据分配了一个线程，一共有 20M 个线程，即 20 * 1024 * 1024 个线程， 这也是 N 的大小。但是我们的一个 block 中只分配了 512 个线程，因此我们需要增加 block 的数量。之前提过，block 的数量是由 grid 的形状控制的，因此你需要在第一个参数中做点修改。你需要多少个 block 呢？ 答案就是 (N + 511) / 512 , 更通用的表达式是这样的：

kernel<<< ( N + block_seize - 1 ) / block_size  , block_size>>>(char* param);

其中 ( N + block_seize - 1 ) / block_size 是一个对 N / block_size 计算向上取整的运算。 所以你应该可以理解为什么上述代码中会用到 (N + 511) / 512 这样奇怪的表达式了。

CUDA程序必须在nvcc的环境下编译运行，并且你需要把.cpp后缀名改为.cu。 建立nvcc的环境需要 Nvidia 的 GPU 。 你可以使用这两行代码再命令行中进行编译和运行。

nvcc test.cu -o test.out    //编译
./test.out                  //运行

上面的代码我直接抄自借鉴 Nvidia 的官方文档，并对其中目前学习未涉及的调用做了删改，以便你能专注于这一种计算方法。在这两章，我们学习了 CUDA 的线程分配、内存调度等基础知识，帮助你编写了第一个kernel程序，当然，我们并不会止步于此，我们的CUDA学习之旅才刚刚开始。

我将从CUDA的运行性能、查询设备参数、处理报错、数据传输、流控制等多个方面继续为你介绍CUDA的入门知识，多方面地带你感受CUDA的 编程之美 。(2024年8月10日)

---

3.测试并行性能指标

主题：性能指标的测定计算 & 事件函数 & 计算带宽和吞吐量

原文链接: How to Implement Performance Metrics in CUDA C/C++

阅读本篇之前，我建议你先安装Nvidia visual-profiler。 使用 nvprof , 你可以在程序运行时跟踪 数据传输、内存使用、API调用 等。 下面是 nvprof 的简介，你可以从文档中找到它的用法。它的具体用法我将不作介绍，希望你在运行程序的时候注意它的功能，利用好这个工具。这对我们分析程序的执行性能指标很有帮助。

• Nvidia visual-profiler简介

nvprof ./test.out

你也许很少在C++程序中一次性运行 20M 的数据，但是这在并行运算中是很常见的。 如何得知这些数据什么时候能够计算完成呢？在C++中，数据按照你的程序执行顺序运算，用有关时间的 API 就可以计算程序运行所花费的时间，例如通过计算时间戳差值获得运行时间。而 CUDA 专门提供了关于线程同步和在事件处打断点的 API， 在CUDA中，使用 event 标记事件点，你也可以像非并行程序那样计算运行时间。

声明两个 event 事件，分别命名为 start, stop，注意它们的类型是 cudaEvent_t， 这是CUDA中专门为 event 定制的类型。然后调用函数 cudaEventCreate() 创建事件。

cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);

接下来要做的，就是将这两个事件标记在你的程序中。 cudaEventRecord() 函数能够把你的事件打在这行代码的断点处。 你可以在你的程序开始和结束的位置打上这两个事件断点，就像下面的代码所展示的一样。这样，你就获得了两个真实记录代码运行时的时刻等信息的事件。

cudaEventRecord(start);
	// ...do something
cudaEventRecord(stop);

下段代码中， cudaEventSynchronize(stop) 在stop事件处做了一次同步，它可以等待stop事件前所有的程序包括并行在 GPU 上的那一部分全部运行完成，调用它可以确保运行到这行代码时你的所有线程都已经运行结束，并不存在仍在进行的线程。然后如何简单而完美地将你的事件运用在测量运行时间呢？就靠下面两行代码。你声明了一个 float 类型的数 milliseconds 来储存这两个事件之间相差的时间，单位时毫秒 ms 。

cudaEventSynchronize(stop);

float milliseconds = 0;
cudaEventElapsedTime(&milliseconds, start, stop);

其实事件的用法和C语言中时间函数的用法类似，但是事件不仅仅储存有这个断点处的时刻信息。如果你感兴趣，你可以尝试更多的 事件API 用法。但是就我们计算带宽和吞吐量来说，获得程序运行前后的时刻信息就已经足够了。

有关带宽 bandwidth 和吞吐量 throughput 的定义我不在这里赘述，请试着在之前的 saxpy 程序中使用这些 API，计算你的矩阵乘加运算程序运行达到的带宽和吞吐量。

如果你足够细心，会发现同步在并行程序中非常重要。如果未对线程同步，将会对程序的逻辑带来混乱。例如，在一个线程中取出了另一个线程还未修改的数据。以下是一些程序同步的API，它们有不同的特性。

deviceSynchronize() 在文档中被称为 heavy hammer ，因为它的同步级别最高，会等待设备上所有线程运行完毕才会开始执行后续的程序。 而 eventSynchronize() 的同步级别没有那么高，例如上面我们使用到的代码，但是它需要传递一个事件参数。而 __synchronize() 是在线程内部进行同步的语句，它的用法后面会介绍，因为涉及到 block 内部 warp 的使用。

现在的你，若需要将线程清晰的划分次序，推荐试一试 eventSynchronize() 。注意，除了线程内部调用的 __synchronize() ，其他的同步函数都使用在 CPU 主程序之中。

现在，让我们改写先前的代码，使用本节讲述的内容，对矩阵乘加并行运算程序计算带宽。

#include <stdio.h>

__global__ void saxpy(int n, float a, float *x, float *y)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n)
        y[i] = a * x[i] + y[i];
}

int main(void)
{
    int N = 20 * (1 << 20);
    float *x, *y, *d_x, *d_y;
    x = (float *)malloc(N * sizeof(float));
    y = (float *)malloc(N * sizeof(float));

    cudaMalloc(&d_x, N * sizeof(float));
    cudaMalloc(&d_y, N * sizeof(float));

    for (int i = 0; i < N; i++)
    {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);

    cudaMemcpy(d_x, x, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, y, N * sizeof(float), cudaMemcpyHostToDevice);

    cudaEventRecord(start);

    saxpy<<<(N + 511) / 512, 512>>>(N, 2.0f, d_x, d_y);

    cudaEventRecord(stop);

    cudaMemcpy(y, d_y, N * sizeof(float), cudaMemcpyDeviceToHost);

    cudaEventSynchronize(stop);
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);

    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
    {
        maxError = max(maxError, abs(y[i] - 4.0f));
    }

    printf("Max error: %fn\n", maxError);
    printf("Effective Bandwidth (GB/s): %fn\n", N * 4 * 3 / milliseconds / 1e6);
}

至于代码中的 N * 4 * 3 / milliseconds / 1e6 表达式，你可以从我引用的官方文档中看到这么计算的解释，也可以去重温一下带宽的定义。回去啃书吧。^_

本节我们利用 event 事件API，计算了程序的带宽，并简单介绍了线程同步的方法和常用API。学习完本节，你就可以开始对你要测算的代码进行一个大比拼，究竟你的并行程序相对于CPU程序快了多少倍呢，值得一试！

在后续的介绍中，我将继续从 查询设备参数、处理报错、数据传输、流控制 等方面为你引入，希望初学者多去研究官方文档，将 CUDA 基础知识和理念熟记于心，并勤加练习。(2024年8月12日)

---

4.探查设备参数、处理报错信息

主题：查询运行设备GPU CUDA运算平台 指定设备

原文链接: How to Query Device Properties and Handle Errors in CUDA C/C++

这一节的内容比较简单，主要介绍了在CUDA中我们如何去查询设备Device的参数，处理运行时的报错。

查询设备参数

为什么要去查询设备参数？这对编写 CUDA 程序有多重要？我的回答是非常重要。CUDA 是面向调用 GPU 运算单元、在 GPU 上开辟线程以运作的。了解我们设备的参数信息，才能知道我们的程序开销，例如调用了多少SM、SM的利用率是多少、你的架构warp的大小是多少。甚至举个更简单的例子，在某些架构中不支持 CUDA 的部分 API，你应该知道如何去规避设备原因导致的不兼容情况。因为我们从顶层的软件、算法，直接深入到了底层硬件，CUDA 作为软硬件之间的桥梁，我们更应该去注意硬件设备的信息。

有关这部分硬件在 CUDA 中对应的调用，你需要首先了解这些硬件在 GPU 中扮演着什么样的角色，以及它们存在的形式。

例如下面这张NVidia Tesla架构微观总览图，简单概括一下，它向你展示了如下信息。

• 拥有7组TPC（Texture/Processor Cluster，纹理处理簇）
• 每个TPC有两组SM（Stream Multiprocessor，流多处理器）
• 每个SM包含：
• • 6个SP（Streaming Processor，流处理器）
• • 2个SFU（Special Function Unit，特殊函数单元）
• • L1缓存、MT Issue（多线程指令获取）、C-Cache（常量缓存）、共享内存
• 除了TPC核心单元，还有与显存、CPU、系统内存交互的各种部件。

还有这份看起来很强大且恐怖的NVidia Kepler架构微观总览图。

非常概括地说，GPU 由显存和计算单元组成。 而 CUDA 中调用的 API 大多不是空中楼阁，它们在实际硬件中有非常明确的对应。了解这些设备信息，有助于你更高效地编写 kernel ，将GPU的并行能力利用得当并发挥到极致。

• 显存（Global Memory）：显存是在GPU板卡上的DRAM，类似于CPU的内存。容量可达16GB，速度慢，CPU和GPU都可以访问
• 计算单元（Streaming Multiprocessor）：执行计算。每一个SM都有自己的控制单元（Control Unit），寄存器（Register），缓存（Cache），指令流水线（execution pipelines）

深入GPU硬件架构及运行机制

我在这里贴上几篇大佬介绍GPU架构的文档，便于去学习。不过，CUDA学习并不要求你对GPU架构的每一块细节都了如指掌，首先建立起宏观的概念和深入挖掘的能力即可。你可以直接复制我的代码运行。

#include <stdio.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <iostream>

using namespace std;

int main()
{
    cudaDeviceProp deviceProp;
    int deviceCount;
    cudaError_t cudaError;
    cudaError = cudaGetDeviceCount(&deviceCount);
    for (int i = 0; i < deviceCount; i++)
    {
        cudaError = cudaGetDeviceProperties(&deviceProp, i);

        cout << "设备 " << i + 1 << " 的主要属性： " << endl;
        cout << "设备显卡型号： " << deviceProp.name << endl;
        cout << "设备上一个线程束（Warp）中包含的线程数量： " << deviceProp.warpSize << endl;
        cout << "多维线程块（Block）数组中，每一维可包含的最大线程（Thread）数量： " << deviceProp.maxThreadsDim[3] << endl;
        cout << "一个线程格（Grid）中，每一维可以包含线程块（Block）数量： " << deviceProp.maxGridSize[3] << endl;
        cout << "设备全局内存总量（以MB为单位）： " << deviceProp.totalGlobalMem / 1024 / 1024 << endl;
        cout << "设备上一个线程块（Block）中可用的最大共享内存（以KB为单位）： " << deviceProp.sharedMemPerBlock / 1024 << endl;
        cout << "设备上一个线程块（Block）种可用的32位寄存器数量： " << deviceProp.regsPerBlock << endl;
        cout << "设备上一个线程块（Block）可包含的最大线程数量： " << deviceProp.maxThreadsPerBlock << endl;
        cout << "设备的计算功能集（Compute Capability）的版本号： " << deviceProp.major << "." << deviceProp.minor << endl;
        cout << "设备上多处理器的数量： " << deviceProp.multiProcessorCount << endl;

        int numSMs;
        cudaDeviceGetAttribute(&numSMs, cudaDevAttrMultiProcessorCount,i);
        cout << "设备上SM的个数： " << numSMs << endl ;
    }

    getchar();
    return 0;
}

处理报错信息

CUDA 的报错信息处理起来很简单，它有一个特定的模板供你使用。这是一个用 inline 注释的内联展开函数，这个模板的功能是把 cuda 的报错信息(cudaError_t 类型)打印出来。如果你调用的 API 并没有出现报错，它会返回 cudaSuccess ,注意这个调用成功的返回值类型也是 cudaError_t 。因此，如果API调用出现错误，以下这个 checkCuda() 的函数会将你的cuda运行结果转换成字符串打印出来，否则就不输出任何信息。

inline cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
    if (result != cudaSuccess)
    {
        fprintf(stderr, "CUDA Runtime Error: %s\n",
                cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
#endif
    return result;
}

你大呼easy，迫不及待想要尝试。只需要用 checkCuda() 接收你调用的API的返回值，一切就水到渠成了。就像这样：

//...
    checkCuda(cudaEventCreate(&startEvent));
    checkCuda(cudaEventCreate(&stopEvent));

    checkCuda(cudaEventRecord(startEvent, 0));
    checkCuda(cudaMemcpy(d, r, bytes, cudaMemcpyHostToDevice));
    checkCuda(cudaEventRecord(stopEvent, 0));
    checkCuda(cudaEventSynchronize(stopEvent));
//...

只要是 cuda 特有的 API ，例如 cudaEventCreate(), cudaMemcpy() …几乎都能用这个方式处理报错信息。大胆尝试一下！

本节我们简单总结了 cuda 中设备信息的查询方式，以及错误信息的处理。跟随文档你能更加细致地了解这些内容，所以请随时别忘记查看文档。文档链接在每一节地标题下都有附上。解下来，我们将从 数据传输、流控制、共享内存、自定义线程组 等方面继续为你总结CUDA 入门学习地知识。(2024年8月14日)

5.更优地数据传输方式

主题：优化数据传输 pageable -> pinned 上锁打包

原文链接: How to Optimize Data Transfers in CUDA C/C++

作为一个内存搬运工，你是否想过怎样更加高效便捷地搬运内存？实际上，不同内存的传输方式具有不同的特点。曾经我们使用 cudaMemcpy() 直接将内存拷贝到 Device, 这样的方法简单且便于理解，但是你需要知道拷贝的过程中数据是如何被传输的。

默认情况下，由 malloc() 创建的 Host 数据是 Pageable (你可以把它想象为未经过打包的内存)。 GPU 无法直接从 pageable 的内存访问数据，因此当调用内存的数据传输时，CUDA 驱动程序必须首先分配一个临时的 pinned 内存(“固定”主机数组, 你可以把它想象为经过打包/上锁/固定的内存)，将主机数据复制到固定数组，然后将数据传输到设备内存，如下图所示。

pinned 内存用作从设备到主机的传输暂存区域。cudaMallocHost() 可以直接创建在 Host 端的 pinned 内存，通过直接创建 pinned 内存，我们可以避免在 pageable 和 pinned 内存之间传输的成本。当你程序运行完毕，应该使用 cudaFreeHost() 解除分配，销毁你创建的 pinned 内存。

创建 pinned 内存可能会失败，因此应始终检查错误，将报错信息打印出来，以便在内存分配错误时及时的发现问题出在了哪里。以下代码演示了 pinned 内存分配和 pageable 内存分配的实现方式，并用 checkCuda() 在每一步调用时打印报错信息。 checkCuda() 的实现在前面小节中已经做了介绍。

以下整理的的代码几乎是初学 CUDA 时所能使用到的所有内存创建方式，不妨将它记在你的笔记当中，以供查阅。

float *h_aPageable, *h_bPageable;    
float *h_aPinned, *h_bPinn
float *d

h_aPageable = (float *)malloc(bytes);                  // Host端创建的pageable内存
h_bPageable = (float *)malloc(bytes);                  // Host端创建的pageable内存
checkCuda(cudaMallocHost((void **)&h_aPinned, bytes)); // Host端创建的 pinned 内存
checkCuda(cudaMallocHost((void **)&h_bPinned, bytes)); // Host端创建的 pinned 内存
checkCuda(cudaMalloc((void **)&d_a, bytes));           // device端创建的内存

运行以下这段代码你可以更加直观地观察到两种内存创建方式对数据传输带宽的影响。pinned 内存在传输到 Device 端的时候并没有像 pageable 内存那样做了一次转换，因此传输效率更高。

#include <stdio.h>
#include <assert.h>

inline cudaError_t checkCuda(cudaError_t result)
{
#if defined(DEBUG) || defined(_DEBUG)
    if (result != cudaSuccess)
    {
        fprintf(stderr, "CUDA Runtime Error: %s\n",
                cudaGetErrorString(result));
        assert(result == cudaSuccess);
    }
#endif
    return result;
}

// 专门定义的拷贝封装
void profileCopies(float *h_a,
                   float *h_b,
                   float *d,
                   unsigned int n,
                   char *desc)
{
    printf("\n%s transfers\n", desc);

    unsigned int bytes = n * sizeof(float);

    // 创建事件
    cudaEvent_t startEvent, stopEvent;

    checkCuda(cudaEventCreate(&startEvent));
    checkCuda(cudaEventCreate(&stopEvent));

    checkCuda(cudaEventRecord(startEvent, 0));
    checkCuda(cudaMemcpy(d, h_a, bytes, cudaMemcpyHostToDevice));
    checkCuda(cudaEventRecord(stopEvent, 0));
    checkCuda(cudaEventSynchronize(stopEvent));

    float time;
    checkCuda(cudaEventElapsedTime(&time, startEvent, stopEvent));
    printf("  Host to Device bandwidth (GB/s): %f\n", bytes * 1e-6 / time);

    checkCuda(cudaEventRecord(startEvent, 0));
    checkCuda(cudaMemcpy(h_b, d, bytes, cudaMemcpyDeviceToHost));
    checkCuda(cudaEventRecord(stopEvent, 0));
    checkCuda(cudaEventSynchronize(stopEvent));

    checkCuda(cudaEventElapsedTime(&time, startEvent, stopEvent));
    printf("  Device to Host bandwidth (GB/s): %f\n", bytes * 1e-6 / time);

    for (int i = 0; i < n; ++i)
    {
        if (h_a[i] != h_b[i])
        {
            printf("*** %s transfers failed ***\n", desc);
            break;
        }
    }

    // 清除事件
    checkCuda(cudaEventDestroy(startEvent));
    checkCuda(cudaEventDestroy(stopEvent));
}

int main()
{
    unsigned int nElements = 4 * 1024 * 1024;
    const unsigned int bytes = nElements * sizeof(float);

    // host 端内存声明
    float *h_aPageable, *h_bPageable;
    float *h_aPinned, *h_bPinned;

    // device 端内存声明
    float *d_a;

    // 分配和初始化内存
    h_aPageable = (float *)malloc(bytes);                  
    h_bPageable = (float *)malloc(bytes);                  
    checkCuda(cudaMallocHost((void **)&h_aPinned, bytes)); 
    checkCuda(cudaMallocHost((void **)&h_bPinned, bytes)); 
    checkCuda(cudaMalloc((void **)&d_a, bytes));           

    for (int i = 0; i < nElements; ++i)
        h_aPageable[i] = i;
    memcpy(h_aPinned, h_aPageable, bytes);
    memset(h_bPageable, 0, bytes);
    memset(h_bPinned, 0, bytes);

    // 记录设备信息 //记录数据传输大小
    cudaDeviceProp prop;
    checkCuda(cudaGetDeviceProperties(&prop, 0));

    printf("\nDevice: %s\n", prop.name);
    printf("Transfer size (MB): %d\n", bytes / (1024 * 1024));

    // 这是一个封装的调用 // 执行拷贝 // 并计算传输时的带宽
    profileCopies(h_aPageable, h_bPageable, d_a, nElements, "Pageable");
    profileCopies(h_aPinned, h_bPinned, d_a, nElements, "Pinned");

    printf("n");

    // 销毁内存好习惯
    cudaFree(d_a);
    cudaFreeHost(h_aPinned);
    cudaFreeHost(h_bPinned);
    free(h_aPageable);
    free(h_bPageable);

    return 0;
}

本节中我们学习了两种创建 Host 内存的 API 的区别，对比了 pinned 和 pageable 两种数据传输方式的性能。但是不应过度分配 pinned 这样的固定内存，这样做会降低整体系统性能，因为它会减少操作系统和其他程序可用的物理内存量。在使用 CUDA 编写程序时，你应该在这两个 API 之间做权衡。(2024年8月19日)

---

欢迎点赞收藏，欢迎访问我的博客，本博客原址点这里。同步更新。