2.1. CUDA C++ 入门

本章通过说明 CUDA 编程模型的一些基本概念如何在 C++ 中呈现，来介绍这些概念。

本编程指南重点介绍 CUDA 运行时 API。CUDA 运行时 API 是在 C++ 中使用 CUDA 最常用的方式，它建立在更底层的 CUDA 驱动程序 API 之上。

CUDA 运行时 API 与 CUDA 驱动程序 API讨论了这两种 API 之间的区别，CUDA 驱动程序 API讨论了如何编写混合使用这两种 API 的代码。

本指南假定已安装 CUDA Toolkit 和 NVIDIA 驱动程序，并且系统中配备了受支持的 NVIDIA GPU。有关安装所需 CUDA 组件的说明，请参阅《CUDA 快速入门指南》。

2.1.1. 使用 NVCC 编译

用 C++ 编写的 GPU 代码使用 NVIDIA Cuda Compiler，即nvcc进行编译。nvcc是一个编译器驱动程序，它简化了编译 C++ 或 PTX 代码的过程：它提供简单且熟悉的命令行选项，并通过调用实现各个编译阶段的一组工具来执行这些选项。

本指南将展示nvcc命令行，这些命令行可用于安装了 CUDA Toolkit 的任何 Linux 系统、Windows 命令行或 power shell，也可用于装有 CUDA Toolkit 的 Windows Subsystem for Linux。本指南的nvcc 章节涵盖了nvcc的常见用例，完整文档由nvcc 用户手册提供。

2.1.2. 核函数

如CUDA 编程模型的介绍中所述，在 GPU 上执行且可从主机调用的函数称为核函数。核函数是为由许多并行线程同时运行而编写的。

2.1.2.1. 指定核函数

核函数的代码使用__global__声明说明符来指定。这向编译器表明，此函数将被编译为在 GPU 上运行，并且能够通过核函数启动来调用。核函数启动是一种开始运行核函数的操作，通常从 CPU 发起。核函数是返回类型为void的函数。

// Kernel definition
__global__ void vecAdd(float* A, float* B, float* C)
{

}

2.1.2.2. 启动核函数

将并行执行核函数的线程数量作为核函数启动的一部分来指定。这称为执行配置。同一核函数的不同调用可以使用不同的执行配置，例如不同数量的线程或线程块。

从 CPU 代码启动核函数有两种方式：三重尖括号记法和cudaLaunchKernelEx。这里介绍三重尖括号记法，它是启动核函数最常用的方式。使用cudaLaunchKernelEx启动核函数的示例在在第 3.1.1 节一节中给出并进行了详细讨论。

2.1.2.2.1. 三重尖括号记法

三重尖括号记法是一种用于启动核函数的CUDA C++ 语言扩展。之所以称为三重尖括号，是因为它使用三个尖括号字符将核函数启动的执行配置括起来，即<<< >>>。执行配置参数在尖括号内以逗号分隔的列表形式指定，类似于函数调用的参数。启动vecAdd核函数的语法如下所示。

 __global__ void vecAdd(float* A, float* B, float* C)
 {

 }

int main()
{
    ...
    // Kernel invocation
    vecAdd<<<1, 256>>>(A, B, C);
    ...
}

三重尖括号记法的前两个参数分别是网格的维度大小和线程块的维度大小。使用一维线程块或网格时，可以用整数指定维度大小。

上述代码启动一个包含 256 个线程的线程块。每个线程都将执行完全相同的核函数代码。在线程和网格索引内建函数中，我们将展示每个线程如何使用其在线程块和网格中的索引来改变它所操作的数据。

每个块的线程数量存在上限，因为一个块的所有线程都驻留在同一个流式多处理器（SM）上，并且必须共享该 SM 的资源。在当前的 GPU 上，一个线程块最多可以包含 1024 个线程。如果资源允许，可以在一个 SM 上同时调度多个线程块。

核函数启动相对于主机线程是异步的。也就是说，核函数将被设置为在 GPU 上执行，但主机代码在继续执行之前，不会等待核函数在 GPU 上执行完成（甚至不会等待其开始执行）。必须使用某种形式的 GPU 与 CPU 之间的同步，才能确定核函数已经完成。最基本的形式是对整个 GPU 进行完全同步，同步 CPU 与 GPU中展示了这种形式。更复杂的同步方法在异步执行中介绍。

使用二维或三维网格或线程块时，使用 CUDA 类型dim3作为网格和线程块的维度大小参数。下面的代码片段展示了MatAdd核函数的启动，使用由 16 × 16 个线程块构成的网格，每个线程块的大小为 8 × 8。

int main()
{
    ...
    dim3 grid(16,16);
    dim3 block(8,8);
    MatAdd<<<grid, block>>>(A, B, C);
    ...
}

2.1.2.3. 线程和网格索引内建函数

在核函数代码内，CUDA 提供内建函数来访问执行配置的参数以及线程或块的索引。

threadIdx给出线程在其线程块内的索引。线程块中的每个线程都有不同的索引。

blockDim给出线程块的维度大小，该大小是在核函数启动的执行配置中指定的。

blockIdx给出线程块在网格内的索引。每个线程块都有不同的索引。

gridDim给出网格的维度大小，该大小是在启动核函数时的执行配置中指定的。

这些内建函数中的每一个都是一个三分量向量，包含.x、.y和.z成员。启动配置中未指定的维度大小默认为 1。threadIdx和blockIdx的索引从零开始。也就是说，threadIdx.x的取值从 0 到blockDim.x-1，包含该上限。.y和.z在各自的维度上以相同方式工作。

类似地，blockIdx.x的取值从 0 到gridDim.x-1，包含该上限；.y和.z维度各自也遵循相同规则。

这些内建函数使单个线程能够确定自己应执行什么工作。回到vecAdd核函数，该核函数接受三个参数，每个参数都是一个浮点数向量。核函数对A和B逐元素相加，并将结果存储到C中。该核函数经过并行化，使每个线程执行一次加法。它计算哪个元素由其线程和网格索引决定。

__global__ void vecAdd(float* A, float* B, float* C)
{
   // calculate which element this thread is responsible for computing
   int workIndex = threadIdx.x + blockDim.x * blockIdx.x;

   // Perform computation
   C[workIndex] = A[workIndex] + B[workIndex];
}

int main()
{
    ...
    // A, B, and C are vectors of 1024 elements
    vecAdd<<<4, 256>>>(A, B, C);
    ...
}

在此示例中，使用 4 个线程块，每个线程块有 256 个线程，来对含有 1024 个元素的向量进行相加。在第一个线程块中，blockIdx.x为零，因此每个线程的 workIndex 就是它的threadIdx.x。在第二个线程块中，blockIdx.x为 1，因此blockDim.x * blockIdx.x与blockDim.x相同，在此情况下为 256。第二个线程块中每个线程的workIndex将是它的threadIdx.x + 256。在第三个线程块中，workIndex将为threadIdx.x + 512。

这种workIndex的计算在一维并行化中十分常见。扩展到二维或三维时，通常会在每个维度上遵循相同的模式。

2.1.2.3.1. 边界检查

上面给出的示例假定向量长度是线程块大小的倍数，此处的线程块大小为 256 个线程。为了使核函数能够处理任意向量长度，我们可以如下所示添加检查，确保内存访问不超出数组边界，然后启动一个包含一些非活动线程的线程块。

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
     // calculate which element this thread is responsible for computing
     int workIndex = threadIdx.x + blockDim.x * blockIdx.x;

     if(workIndex < vectorLength)
     {
         // Perform computation
         C[workIndex] = A[workIndex] + B[workIndex];
     }
}

使用上述核函数代码，可以启动多于所需数量的线程，而不会造成对数组的越界访问。当workIndex超过vectorLength时，线程会退出而不执行任何工作。在线程块中启动不执行工作的额外线程不会产生很大的开销，但应避免启动其中所有线程都不执行工作的线程块。现在，这个核函数可以处理长度不是块大小倍数的向量。

所需线程块的数量可以通过将所需线程数（此处为向量长度）除以每个块的线程数，然后向上取整来计算。也就是说，将所需线程数除以每个块的线程数进行整数除法，并向上取整。下面给出了一种将其表示为单次整数除法的常见方式。在整数除法之前加上threads - 1，就能使其行为类似于向上取整函数，仅当向量长度不能被每个块的线程数整除时，才增加一个线程块。

// vectorLength is an integer storing number of elements in the vector
int threads = 256;
int blocks = (vectorLength + threads-1)/threads;
vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);

CUDA Core Compute Library (CCCL)提供了一个方便的实用工具cuda::ceil_div，用于执行这种向上取整除法，以计算核函数启动所需的块数。包含头文件<cuda/cmath>即可使用此实用工具。

// vectorLength is an integer storing number of elements in the vector
int threads = 256;
int blocks = cuda::ceil_div(vectorLength, threads);
vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);

这里为每个块选择 256 个线程是任意的，不过这个数值通常很适合作为初始选择。

2.1.3. GPU 计算中的内存

要使用上面展示的vecAdd核函数，数组A、B和C必须位于 GPU 可访问的内存中。有几种不同的方法可以做到这一点，这里将介绍其中两种。其他方法将在后面有关统一内存的章节中介绍。在 GPU 上运行的代码可使用的内存空间已在GPU 内存中介绍，并将在GPU 设备内存空间中更详细地讨论。

2.1.3.1. 统一内存

统一内存是 CUDA 运行时的一项功能，它让 NVIDIA 驱动程序管理主机与一个或多个设备之间的数据移动。内存通过cudaMallocManagedAPI 分配，或通过使用__managed__说明符声明变量来分配。无论 GPU 还是 CPU 尝试访问该内存，NVIDIA 驱动程序都会确保它能够访问该内存。

下面的代码展示了一个用于启动vecAdd核函数的完整函数，它为将在 GPU 上使用的输入和输出向量使用统一内存。cudaMallocManaged分配可从 CPU 或 GPU 访问的缓冲区。这些缓冲区使用cudaFree释放。

void unifiedMemExample(int vectorLength)
{
    // Pointers to memory vectors
    float* A = nullptr;
    float* B = nullptr;
    float* C = nullptr;
    float* comparisonResult = (float*)malloc(vectorLength*sizeof(float));

    // Use unified memory to allocate buffers
    cudaMallocManaged(&A, vectorLength*sizeof(float));
    cudaMallocManaged(&B, vectorLength*sizeof(float));
    cudaMallocManaged(&C, vectorLength*sizeof(float));

    // Initialize vectors on the host
    initArray(A, vectorLength);
    initArray(B, vectorLength);

    // Launch the kernel. Unified memory will make sure A, B, and C are
    // accessible to the GPU
    int threads = 256;
    int blocks = cuda::ceil_div(vectorLength, threads);
    vecAdd<<<blocks, threads>>>(A, B, C, vectorLength);
    // Wait for the kernel to complete execution
    cudaDeviceSynchronize();

    // Perform computation serially on CPU for comparison
    serialVecAdd(A, B, comparisonResult, vectorLength);

    // Confirm that CPU and GPU got the same answer
    if(vectorApproximatelyEqual(C, comparisonResult, vectorLength))
    {
        printf("Unified Memory: CPU and GPU answers match\n");
    }
    else
    {
        printf("Unified Memory: Error - CPU and GPU answers do not match\n");
    }

    // Clean Up
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    free(comparisonResult);

}

CUDA 支持的所有操作系统和 GPU 都支持统一内存，不过其底层机制和性能可能因系统架构而异。统一内存提供了更多详细信息。在某些 Linux 系统上（例如具有地址转换服务或异构内存管理的系统），所有系统内存都会自动成为统一内存，无需使用cudaMallocManaged或__managed__说明符。

2.1.3.2. 显式内存管理

显式管理内存分配和内存空间之间的数据迁移有助于提高应用程序性能，不过这确实会使代码更加冗长。下面的代码使用cudaMalloc在 GPU 上显式分配内存。GPU 上的内存使用cudaFreeAPI 释放，这与前一个示例中用于统一内存的 API 相同。

void explicitMemExample(int vectorLength)
{
    // Pointers for host memory
    float* A = nullptr;
    float* B = nullptr;
    float* C = nullptr;
    float* comparisonResult = (float*)malloc(vectorLength*sizeof(float));
    
    // Pointers for device memory
    float* devA = nullptr;
    float* devB = nullptr;
    float* devC = nullptr;

    //Allocate Host Memory using cudaMallocHost API. This is best practice
    // when buffers will be used for copies between CPU and GPU memory
    cudaMallocHost(&A, vectorLength*sizeof(float));
    cudaMallocHost(&B, vectorLength*sizeof(float));
    cudaMallocHost(&C, vectorLength*sizeof(float));

    // Initialize vectors on the host
    initArray(A, vectorLength);
    initArray(B, vectorLength);

    // start-allocate-and-copy
    // Allocate memory on the GPU
    cudaMalloc(&devA, vectorLength*sizeof(float));
    cudaMalloc(&devB, vectorLength*sizeof(float));
    cudaMalloc(&devC, vectorLength*sizeof(float));

    // Copy data to the GPU
    cudaMemcpy(devA, A, vectorLength*sizeof(float), cudaMemcpyDefault);
    cudaMemcpy(devB, B, vectorLength*sizeof(float), cudaMemcpyDefault);
    cudaMemset(devC, 0, vectorLength*sizeof(float));
    // end-allocate-and-copy

    // Launch the kernel
    int threads = 256;
    int blocks = cuda::ceil_div(vectorLength, threads);
    vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);
    // wait for kernel execution to complete
    cudaDeviceSynchronize();

    // Copy results back to host
    cudaMemcpy(C, devC, vectorLength*sizeof(float), cudaMemcpyDefault);

    // Perform computation serially on CPU for comparison
    serialVecAdd(A, B, comparisonResult, vectorLength);

    // Confirm that CPU and GPU got the same answer
    if(vectorApproximatelyEqual(C, comparisonResult, vectorLength))
    {
        printf("Explicit Memory: CPU and GPU answers match\n");
    }
    else
    {
        printf("Explicit Memory: Error - CPU and GPU answers to not match\n");
    }

    // clean up
    cudaFree(devA);
    cudaFree(devB);
    cudaFree(devC);
    cudaFreeHost(A);
    cudaFreeHost(B);
    cudaFreeHost(C);
    free(comparisonResult);
}

CUDA APIcudaMemcpy用于将数据从位于 CPU 上的缓冲区复制到位于 GPU 上的缓冲区。除了目标指针、源指针和以字节为单位的大小之外，cudaMemcpy的最后一个参数是cudaMemcpyKind_t。它可以具有如下取值：

cudaMemcpyHostToDevice用于从 CPU 到 GPU 的复制
cudaMemcpyDeviceToHost用于从 GPU 到 CPU 的复制
cudaMemcpyDeviceToDevice用于 GPU 内部或 GPU 之间的复制

在此示例中，cudaMemcpyDefault作为最后一个参数传递给cudaMemcpy。这使 CUDA 使用源指针和目标指针的值来确定要执行的复制类型。

cudaMemcpyAPI 是同步的。也就是说，它直到复制完成后才会返回。异步复制在在 CUDA 流中启动内存传输中介绍。

该代码使用cudaMallocHost在 CPU 上分配内存。这会在主机上分配页锁定内存，这种内存可以提高复制性能，并且是异步内存传输所必需的。通常，对于将用于与 GPU 进行双向数据传输的 CPU 缓冲区，使用页锁定内存是一种良好做法。如果过多的主机内存被页锁定，某些系统上的性能可能会下降。最佳做法是仅对将用于向 GPU 发送数据或从 GPU 接收数据的缓冲区进行页锁定。

2.1.3.3. 内存管理与应用程序性能

从上面的示例可以看出，显式内存管理更加繁琐，要求程序员指定主机与设备之间的复制操作。这既是显式内存管理的优势，也是其劣势：它能够更精确地控制何时在主机与设备之间复制数据、内存驻留在何处，以及具体在哪个位置分配哪些内存。显式内存管理可以通过控制内存传输以及使其与其他计算重叠，提供提升性能的机会。

使用统一内存时，可以使用 CUDA API（将在内存建议与预取中介绍）向负责管理内存的 NVIDIA 驱动程序提供提示，从而在使用统一内存时获得显式内存管理所带来的部分性能优势。

2.1.4. 同步 CPU 与 GPU

如启动核函数中所述，核函数启动相对于发起调用的 CPU 线程是异步的。这意味着，在核函数完成之前，甚至可能在其启动之前，CPU 线程的控制流就会继续执行。要确保在主机代码继续执行之前核函数已经执行完毕，就需要某种同步机制。

同步 GPU 与主机线程的最简单方法是使用cudaDeviceSynchronize，它会阻塞主机线程，直到此前向 GPU 提交的所有工作都已完成。对于本章中的示例，这样做已经足够，因为 GPU 上仅执行单个操作。在较大的应用程序中，可能有多个流在 GPU 上执行工作，而cudaDeviceSynchronize会等待所有流中的工作完成。在这些应用程序中，建议使用流同步 API 仅与特定流同步，或者使用CUDA 事件。这些内容将在异步执行一章中详细介绍。

2.1.5. 整合所有内容

下面的代码清单给出了本章介绍的简单向量加法核函数的完整代码，以及所有主机代码和用于检查所获结果是否正确的实用函数。这些示例默认使用 1024 的向量长度，但也接受通过可执行文件的命令行参数指定不同的向量长度。

统一内存

#include <cuda_runtime_api.h>
#include <memory.h>
#include <cstdlib>
#include <ctime>
#include <stdio.h>
#include <cuda/cmath>

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
    int workIndex = threadIdx.x + blockIdx.x*blockDim.x;
    if(workIndex < vectorLength)
    {
        C[workIndex] = A[workIndex] + B[workIndex];
    }
}

void initArray(float* A, int length)
{
     std::srand(std::time({}));
    for(int i=0; i<length; i++)
    {
        A[i] = rand() / (float)RAND_MAX;
    }
}

void serialVecAdd(float* A, float* B, float* C,  int length)
{
    for(int i=0; i<length; i++)
    {
        C[i] = A[i] + B[i];
    }
}

bool vectorApproximatelyEqual(float* A, float* B, int length, float epsilon=0.00001)
{
    for(int i=0; i<length; i++)
    {
        if(fabs(A[i] -B[i]) > epsilon)
        {
            printf("Index %d mismatch: %f != %f", i, A[i], B[i]);
            return false;
        }
    }
    return true;
}

//unified-memory-begin
void unifiedMemExample(int vectorLength)
{
    // Pointers to memory vectors
    float* A = nullptr;
    float* B = nullptr;
    float* C = nullptr;
    float* comparisonResult = (float*)malloc(vectorLength*sizeof(float));

    // Use unified memory to allocate buffers
    cudaMallocManaged(&A, vectorLength*sizeof(float));
    cudaMallocManaged(&B, vectorLength*sizeof(float));
    cudaMallocManaged(&C, vectorLength*sizeof(float));

    // Initialize vectors on the host
    initArray(A, vectorLength);
    initArray(B, vectorLength);

    // Launch the kernel. Unified memory will make sure A, B, and C are
    // accessible to the GPU
    int threads = 256;
    int blocks = cuda::ceil_div(vectorLength, threads);
    vecAdd<<<blocks, threads>>>(A, B, C, vectorLength);
    // Wait for the kernel to complete execution
    cudaDeviceSynchronize();

    // Perform computation serially on CPU for comparison
    serialVecAdd(A, B, comparisonResult, vectorLength);

    // Confirm that CPU and GPU got the same answer
    if(vectorApproximatelyEqual(C, comparisonResult, vectorLength))
    {
        printf("Unified Memory: CPU and GPU answers match\n");
    }
    else
    {
        printf("Unified Memory: Error - CPU and GPU answers do not match\n");
    }

    // Clean Up
    cudaFree(A);
    cudaFree(B);
    cudaFree(C);
    free(comparisonResult);

}
//unified-memory-end


int main(int argc, char** argv)
{
    int vectorLength = 1024;
    if(argc >=2)
    {
        vectorLength = std::atoi(argv[1]);
    }
    unifiedMemExample(vectorLength);		
    return 0;
}

显式内存管理

#include <cuda_runtime_api.h>
#include <memory.h>
#include <cstdlib>
#include <ctime>
#include <stdio.h>
#include <cuda/cmath>

__global__ void vecAdd(float* A, float* B, float* C, int vectorLength)
{
    int workIndex = threadIdx.x + blockIdx.x*blockDim.x;
    if(workIndex < vectorLength)
    {
        C[workIndex] = A[workIndex] + B[workIndex];
    }
}

void initArray(float* A, int length)
{
     std::srand(std::time({}));
    for(int i=0; i<length; i++)
    {
        A[i] = rand() / (float)RAND_MAX;
    }
}

void serialVecAdd(float* A, float* B, float* C,  int length)
{
    for(int i=0; i<length; i++)
    {
        C[i] = A[i] + B[i];
    }
}

bool vectorApproximatelyEqual(float* A, float* B, int length, float epsilon=0.00001)
{
    for(int i=0; i<length; i++)
    {
        if(fabs(A[i] -B[i]) > epsilon)
        {
            printf("Index %d mismatch: %f != %f", i, A[i], B[i]);
            return false;
        }
    }
    return true;
}

//explicit-memory-begin
void explicitMemExample(int vectorLength)
{
    // Pointers for host memory
    float* A = nullptr;
    float* B = nullptr;
    float* C = nullptr;
    float* comparisonResult = (float*)malloc(vectorLength*sizeof(float));
    
    // Pointers for device memory
    float* devA = nullptr;
    float* devB = nullptr;
    float* devC = nullptr;

    //Allocate Host Memory using cudaMallocHost API. This is best practice
    // when buffers will be used for copies between CPU and GPU memory
    cudaMallocHost(&A, vectorLength*sizeof(float));
    cudaMallocHost(&B, vectorLength*sizeof(float));
    cudaMallocHost(&C, vectorLength*sizeof(float));

    // Initialize vectors on the host
    initArray(A, vectorLength);
    initArray(B, vectorLength);

    // start-allocate-and-copy
    // Allocate memory on the GPU
    cudaMalloc(&devA, vectorLength*sizeof(float));
    cudaMalloc(&devB, vectorLength*sizeof(float));
    cudaMalloc(&devC, vectorLength*sizeof(float));

    // Copy data to the GPU
    cudaMemcpy(devA, A, vectorLength*sizeof(float), cudaMemcpyDefault);
    cudaMemcpy(devB, B, vectorLength*sizeof(float), cudaMemcpyDefault);
    cudaMemset(devC, 0, vectorLength*sizeof(float));
    // end-allocate-and-copy

    // Launch the kernel
    int threads = 256;
    int blocks = cuda::ceil_div(vectorLength, threads);
    vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);
    // wait for kernel execution to complete
    cudaDeviceSynchronize();

    // Copy results back to host
    cudaMemcpy(C, devC, vectorLength*sizeof(float), cudaMemcpyDefault);

    // Perform computation serially on CPU for comparison
    serialVecAdd(A, B, comparisonResult, vectorLength);

    // Confirm that CPU and GPU got the same answer
    if(vectorApproximatelyEqual(C, comparisonResult, vectorLength))
    {
        printf("Explicit Memory: CPU and GPU answers match\n");
    }
    else
    {
        printf("Explicit Memory: Error - CPU and GPU answers to not match\n");
    }

    // clean up
    cudaFree(devA);
    cudaFree(devB);
    cudaFree(devC);
    cudaFreeHost(A);
    cudaFreeHost(B);
    cudaFreeHost(C);
    free(comparisonResult);
}
//explicit-memory-end


int main(int argc, char** argv)
{
    int vectorLength = 1024;
    if(argc >=2)
    {
        vectorLength = std::atoi(argv[1]);
    }
    explicitMemExample(vectorLength);		
    return 0;
}

可以使用nvcc构建并运行这些示例，如下所示：

$ nvcc vecAdd_unifiedMemory.cu -o vecAdd_unifiedMemory
$ ./vecAdd_unifiedMemory
Unified Memory: CPU and GPU answers match
$ ./vecAdd_unifiedMemory 4096
Unified Memory: CPU and GPU answers match

$ nvcc vecAdd_explicitMemory.cu -o vecAdd_explicitMemory
$ ./vecAdd_explicitMemory
Explicit Memory: CPU and GPU answers match
$ ./vecAdd_explicitMemory 4096
Explicit Memory: CPU and GPU answers match

在这些示例中，所有线程都在执行独立的工作，不需要相互协调或同步。线程经常需要与其他线程协作和通信来完成工作。一个线程块内的线程可以通过共享内存共享数据，并通过同步来协调内存访问。

线程块级别最基本的同步机制是__syncthreads()内建函数，它充当一个屏障，线程块中的所有线程都必须在此等待，之后才允许任何线程继续执行。共享内存给出了使用共享内存的示例。

为了高效协作，共享内存应当是靠近各处理器核心的低延迟内存（很像 L1 缓存），而__syncthreads()应当是轻量级的。__syncthreads()仅同步单个线程块内的线程。

仅在特定情况下才支持线程块之间的同步。例如，线程块簇允许簇内的线程块进行同步，而协作组 API提供了创建跨线程块同步域的机制。

通常，将同步限制在线程块内部能够获得最佳性能。线程块仍然可以使用原子内存函数共同处理结果，这些函数将在后续各节中介绍。

在第 3.2.4 节这一节中，介绍了 CUDA 同步原语，它们提供非常细粒度的控制，以最大限度地提高性能和资源利用率。

2.1.6. 运行时初始化

CUDA 运行时为系统中的每个设备创建一个CUDA 上下文。该上下文是此设备的主上下文，在首次调用需要此设备上的活动上下文的运行时函数时进行初始化。该上下文由应用程序的所有主机线程共享。作为上下文创建过程的一部分，设备代码会在必要时进行即时编译，并加载到设备内存中。所有这一切都是透明进行的。可以通过驱动程序 API 访问 CUDA 运行时创建的主上下文，以实现互操作，如运行时 API 与驱动程序 API 之间的互操作性中所述。

从 CUDA 12.0 开始，cudaInitDevice和cudaSetDevice调用会初始化运行时以及与指定设备关联的主上下文。如果运行时 API 请求发生在这些调用之前，运行时将隐式使用设备 0，并根据需要自行初始化以处理这些请求。在对运行时函数调用计时以及解读首次调用运行时所返回的错误码时，这一点很重要。在 CUDA 12.0 之前，cudaSetDevice不会初始化运行时。

cudaDeviceReset会销毁当前设备的主上下文。如果在主上下文被销毁后调用 CUDA 运行时 API，则会为该设备创建一个新的主上下文。

注意

CUDA 接口使用全局状态，该状态在主机程序启动期间初始化，并在主机程序终止期间销毁。在程序启动期间或 main 之后的程序终止期间使用任何这些接口（无论隐式还是显式），都会导致未定义行为。

从 CUDA 12.0 开始，cudaSetDevice在更改主机线程的当前设备后，如果运行时尚未初始化，就会显式初始化运行时。早期版本的 CUDA 会延迟新设备上的运行时初始化，直到cudaSetDevice之后首次调用运行时函数时才进行。因此，检查cudaSetDevice的返回值以发现初始化错误非常重要。

参考手册中错误处理和版本管理部分的运行时函数不会初始化运行时。

2.1.7. CUDA 中的错误检查

每个 CUDA API 都会返回一个枚举类型cudaError_t的值。在示例代码中，通常不会检查这些错误。在生产应用程序中，最佳实践是始终检查和处理每次 CUDA API 调用的返回值。没有错误时，返回值为cudaSuccess。许多应用程序选择实现一个实用宏，如下所示

#define CUDA_CHECK(expr_to_check) do {            \
    cudaError_t result  = expr_to_check;          \
    if(result != cudaSuccess)                     \
    {                                             \
        fprintf(stderr,                           \
                "CUDA Runtime Error: %s:%i:%d = %s\n", \
                __FILE__,                         \
                __LINE__,                         \
                result,\
                cudaGetErrorString(result));      \
    }                                             \
} while(0)

此宏使用cudaGetErrorString API，该 API 会返回一个便于人阅读的字符串，用于描述特定cudaError_t值的含义。使用上述宏时，应用程序会将 CUDA 运行时 API 调用放在CUDA_CHECK(expression)宏中，如下所示：

    CUDA_CHECK(cudaMalloc(&devA, vectorLength*sizeof(float)));
    CUDA_CHECK(cudaMalloc(&devB, vectorLength*sizeof(float)));
    CUDA_CHECK(cudaMalloc(&devC, vectorLength*sizeof(float)));

如果这些调用中的任何一个检测到错误，就会使用此宏将错误打印到stderr。此宏在较小的项目中很常见，但也可以对其进行调整，以适配较大应用程序中的日志系统或其他错误处理机制。

注意

需要特别注意，任何 CUDA API 调用返回的错误状态，也可能表明先前提交的异步操作发生了错误。异步错误处理一节对此有更详细的介绍。

2.1.7.1. 错误状态

CUDA 运行时为每个主机线程维护一个cudaError_t状态。该值默认为cudaSuccess，每当发生错误时都会被覆盖。cudaGetLastError返回当前错误状态，然后将其重置为cudaSuccess。另一种方式是使用cudaPeekAtLastError，它返回错误状态但不将其重置。

使用三尖括号语法进行核函数启动时，不会返回cudaError_t。良好的实践是在核函数启动后立即检查错误状态，以检测核函数启动中的即时错误，或核函数启动之前的异步错误。在核函数启动后立即检查错误状态时，值为cudaSuccess并不意味着核函数已成功执行，甚至不意味着它已经开始执行。这只能确认传递给运行时的核函数启动参数和执行配置没有触发任何错误，并且错误状态不是核函数开始之前的先前错误或异步错误。

2.1.7.2. 异步错误

CUDA 核函数启动和许多运行时 API 都是异步的。异步 CUDA 运行时 API 将在异步执行中详细讨论。每当发生错误时，CUDA 错误状态都会被设置并覆盖。这意味着，异步操作执行期间发生的错误，只有在下一次检查错误状态时才会报告。如前所述，这可能是对cudaGetLastError、cudaPeekAtLastError的调用，也可能是任何返回cudaError_t的 CUDA API 调用。

CUDA 运行时 API 函数返回错误时，不会清除错误状态。这意味着，对于异步错误（例如核函数的无效内存访问），每个 CUDA 运行时 API 都会返回其错误码，直到通过调用cudaGetLastError清除错误状态为止。

    vecAdd<<<blocks, threads>>>(devA, devB, devC);
    // check error state after kernel launch
    CUDA_CHECK(cudaGetLastError());
    // wait for kernel execution to complete
    // The CUDA_CHECK will report errors that occurred during execution of the kernel
    CUDA_CHECK(cudaDeviceSynchronize());
    

注意

cudaStreamQuery和cudaEventQuery可能返回的cudaError_t值cudaErrorNotReady不被视为错误，cudaPeekAtLastError或cudaGetLastError也不会报告该值。

2.1.7.3. `CUDA_LOG_FILE`

另一种识别 CUDA 错误的好方法是使用CUDA_LOG_FILE环境变量。设置此环境变量后，CUDA 驱动程序会将遇到的错误消息写入文件，文件路径由该环境变量指定。例如，考虑下面这段不正确的 CUDA 代码，它试图启动一个线程块，其大小超过了所有架构所支持的最大值。

__global__ void k()
{ }

int main()
{
        k<<<8192, 4096>>>(); // Invalid block size
        CUDA_CHECK(cudaGetLastError());
        return 0;
}

构建并运行此代码后，核函数启动后的检查会使用第 2.1.7 节中展示的宏检测并报告错误。

$ nvcc errorLogIllustration.cu -o errlog
$ ./errlog
CUDA Runtime Error: /home/cuda/intro-cpp/errorLogIllustration.cu:24:1 = invalid argument

不过，在将CUDA_LOG_FILE设置为一个文本文件后运行应用程序，该文件会包含关于错误的更多一些信息。

$ env CUDA_LOG_FILE=cudaLog.txt ./errlog
CUDA Runtime Error: /home/cuda/intro-cpp/errorLogIllustration.cu:24:1 = invalid argument
$ cat cudaLog.txt
[12:46:23.854][137216133754880][CUDA][E] One or more of block dimensions of (4096,1,1) exceeds corresponding maximum value of (1024,1024,64)
[12:46:23.854][137216133754880][CUDA][E] Returning 1 (CUDA_ERROR_INVALID_VALUE) from cuLaunchKernel

将CUDA_LOG_FILE设置为stdout或stderr，会分别打印到标准输出和标准错误。使用CUDA_LOG_FILE环境变量，即使应用程序没有针对 CUDA 返回值实现适当的错误检查，也能够捕获并识别 CUDA 错误。这种方法对于调试可能极为有效，但仅靠该环境变量并不能让应用程序在运行时处理 CUDA 错误并从中恢复。CUDA 的错误日志管理功能还允许向驱动程序注册一个回调函数，每当检测到错误时就会调用该函数。这可用于在运行时捕获和处理错误，还可将 CUDA 错误日志记录无缝集成到应用程序现有的日志系统中。

第 4.9 节给出了 CUDA 错误日志管理功能的更多示例。错误日志管理和CUDA_LOG_FILE可用于 NVIDIA 驱动程序 r570 及更高版本。

2.1.8. 设备函数与主机函数

__global__说明符用于指示核函数的入口点，即将被调用以在 GPU 上并行执行的函数。核函数通常从主机启动，不过，也可以使用动态并行在另一个核函数内部启动核函数。

说明符__device__表示一个函数应当编译为在 GPU 上运行，并且可由其他__device__或__global__函数调用。函数，包括类成员函数、函数对象和 lambda 表达式，可以同时指定为__device__和__host__，如下例所示。

2.1.9. 变量说明符

可以在静态变量声明中使用CUDA 说明符来控制存储位置。

__device__指定变量存储在全局内存中
__constant__指定变量存储在常量内存中
__managed__指定变量以统一内存形式存储
__shared__指定变量存储在共享内存中

在__device__或__global__函数内部声明变量且不带说明符时，会尽可能将其分配到寄存器中，并在必要时分配到局部内存中。在__device__或__global__函数外部声明的任何不带说明符的变量，都将分配在系统内存中。

2.1.9.1. 检测设备编译

为函数指定__host__ __device__时，就是指示编译器为此函数同时生成 GPU 代码和 CPU 代码。在这样的函数中，可能希望使用预处理器，为函数的 GPU 版本或 CPU 版本单独指定代码。检查是否定义了__CUDA_ARCH_是最常见的方法，如下面的示例所示。

2.1.10. 线程块簇

从计算能力 9.0 开始，CUDA 编程模型包含一个可选的层级，称为线程块簇，由线程块组成。正如线程块内的线程保证会被共同调度到一个流式多处理器（SM）上一样，簇内的线程块也保证会被共同调度到 GPU 中的一个 GPU 处理簇（GPC）上。

与线程块类似，簇也被组织为一维、二维或三维的线程块簇网格，如图 5所示。

簇中的线程块数量可以由用户定义，CUDA 支持的可移植簇大小为每个簇最多 8 个线程块。请注意，对于规模过小、无法支持 8 个多处理器的 GPU 硬件或 MIG 配置，最大簇大小会相应减小。如何识别这些较小的配置，以及支持线程块簇大小超过 8 的较大配置，取决于具体架构，可以使用cudaOccupancyMaxPotentialClusterSize API 进行查询。

簇中的所有线程块保证会被共同调度到单个 GPU 处理簇（GPC）上同时执行，并允许簇内的线程块使用协作组 APIcluster.sync()执行硬件支持的同步。簇组还提供成员函数，可分别使用num_threads()和num_blocks() API，以线程数或线程块数查询簇组大小。可以分别使用dim_threads()和dim_blocks() API 查询线程或线程块在簇组中的秩。

属于同一个簇的线程块可以访问分布式共享内存，它由簇内所有线程块的共享内存组合而成。簇中的线程块能够对分布式共享内存中的任意地址执行读、写和原子操作。分布式共享内存给出了在分布式共享内存中计算直方图的示例。

注意

在使用簇支持启动的核函数中，出于兼容性考虑，gridDim 变量仍然以线程块数量表示大小。可以使用协作组 API 获取一个块在簇中的秩。

2.1.10.1. 使用三重尖括号表示法进行带簇的启动

可以通过使用__cluster_dims__(X,Y,Z)的编译时核函数属性，或通过 CUDA 核函数启动 APIcudaLaunchKernelEx，在核函数中启用线程块簇。下面的示例展示了如何使用编译时核函数属性启动簇。使用核函数属性设置的簇大小在编译时就已固定，之后可以使用传统的<<< , >>>启动核函数。如果核函数使用编译时簇大小，则在启动该核函数时无法修改簇大小。

// Kernel definition
// Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{

}

int main()
{
    float *input, *output;
    // Kernel invocation with compile time cluster size
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // The grid dimension is not affected by cluster launch, and is still enumerated
    // using number of blocks.
    // The grid dimension must be a multiple of cluster size.
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

2.1. CUDA C++ 入门

2.1.1. 使用 NVCC 编译

2.1.2. 核函数

2.1.2.1. 指定核函数

2.1.2.2. 启动核函数

2.1.2.2.1. 三重尖括号记法

2.1.2.3. 线程和网格索引内建函数

2.1.2.3.1. 边界检查

2.1.3. GPU 计算中的内存

2.1.3.1. 统一内存

2.1.3.2. 显式内存管理

2.1.3.3. 内存管理与应用程序性能

2.1.4. 同步 CPU 与 GPU

2.1.5. 整合所有内容

2.1.6. 运行时初始化

2.1.7. CUDA 中的错误检查

2.1.7.1. 错误状态

2.1.7.2. 异步错误

2.1.7.3. CUDA_LOG_FILE

2.1.8. 设备函数与主机函数

2.1.9. 变量说明符

2.1.9.1. 检测设备编译

2.1.10. 线程块簇

2.1.10.1. 使用三重尖括号表示法进行带簇的启动

2.1.7.3. `CUDA_LOG_FILE`