CUDA-编程学习手册-二-

CUDA 编程学习手册（二）

原文：annas-archive.org/md5/f6da79e769f988319eb178273ecbf55b

译者：飞龙

协议：CC BY-NC-SA 4.0

第四章：核心执行模型和优化策略

CUDA 编程有一个主机操作的过程。例如，我们需要分配全局内存，将数据传输到 GPU，执行核心函数，将数据传输回主机，清理全局内存。这是因为 GPU 是系统中的一个额外处理单元，所以我们需要关心它的执行和数据传输。这是与 CPU 编程相比另一个不同的 GPU 编程方面。

在本章中，我们将涵盖 CUDA 核心执行模型和 CUDA 流，它们控制 CUDA 操作。然后，我们将讨论系统级别的优化策略。接下来，我们将涵盖 CUDA 事件来测量 GPU 事件时间，以及如何使用 CUDA 事件来测量核心执行时间。之后，我们将涵盖各种 CUDA 核心执行模型，并讨论这些特性对 GPU 操作的影响。

本章将涵盖以下主题：

• 使用 CUDA 流的核心执行

• 流水线化 GPU 执行

• CUDA 回调函数

• 具有优先级的 CUDA 流

• 使用 CUDA 事件估计核心执行时间

• CUDA 动态并行性

• 网格级协作组

• 使用 OpenMP 的 CUDA 核心调用

• 多进程服务

• 核心执行开销比较

技术要求

本章要求我们使用的 CUDA 版本应该晚于 9.x，并且 GPU 架构应该是 Volta 或 Turing。如果你使用的是 Pascal 架构的 GPU，那么跳过Grid-level cooperative groups部分，因为这个特性是为 Volta 架构引入的。

使用 CUDA 流的核心执行

在 CUDA 编程中，流是与 GPU 相关的一系列命令。换句话说，所有的核心调用和数据传输都由 CUDA 流处理。默认情况下，CUDA 提供了一个默认流，所有的命令都隐式地使用这个流。因此，我们不需要自己处理这个。

CUDA 支持显式创建额外的流。虽然流中的操作是顺序的，但 CUDA 可以通过使用多个流同时执行多个操作。让我们学习如何处理流，以及它们具有哪些特性。

CUDA 流的使用

以下代码展示了如何创建、使用和终止 CUDA 流的示例：

cudaStream_t stream;
cudaStreamCreate(&stream);
foo_kernel<<< grid_size, block_size, 0, stream >>>();
cudaStreamDestroy(stream);

正如你所看到的，我们可以使用cudaStream_t来处理 CUDA 流。而且，我们可以使用cudaStreamCreate()来创建它，使用cudaStreamDestroy()来终止它。注意我们应该提供一个指向cudaStreamCreate()的指针。创建的流会传递给核心函数的第四个参数。

然而，我们之前并没有提供这样的流。这是因为 CUDA 提供了一个默认流，以便所有的 CUDA 操作都可以进行。现在，让我们编写一个使用默认流和多个流的应用程序。然后，我们将看到我们的应用程序如何改变。

首先，让我们编写一个使用默认 CUDA 流的应用程序，如下所示：

__global__ void foo_kernel(int step)
{
    printf("loop: %d\n", step);
}

int main()
{
    for (int i = 0; i < 5; i++)
 // CUDA kernel call with the default stream
 foo_kernel<<< 1, 1, 0, 0 >>>(i);
    cudaDeviceSynchronize();
    return 0;
}

正如你在代码中看到的，我们以流 ID 为0调用了核心函数，因为默认流的标识值为0。编译代码并查看执行输出：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_default_stream ./1_cuda_default_stream.cu

输出是什么？我们可以期待输出将是循环索引的顺序。以下时间轴视图显示了这段代码的操作：

可以预期，在同一个流中进行循环操作将显示核心执行的顺序。那么，如果我们使用多个 CUDA 流，并且每个循环步骤使用不同的流，会有什么改变？以下代码展示了使用不同流从 CUDA 核心函数打印循环索引的示例：

__global__ void foo_kernel(int step)
{
    printf("loop: %d\n", step);
}

int main()
{
    int n_stream = 5;
    cudaStream_t *ls_stream;
    ls_stream = (cudaStream_t*) new cudaStream_t[n_stream];

    // create multiple streams
    for (int i = 0; i < n_stream; i++)
        cudaStreamCreate(&ls_stream[i]);

    // execute kernels with the CUDA stream each
    for (int i = 0; i < n_stream; i++)
        foo_kernel<<< 1, 1, 0, ls_stream[i] >>>(i);

    // synchronize the host and GPU
    cudaDeviceSynchronize();

    // terminates all the created CUDA streams
    for (int i = 0; i < n_stream; i++)
        cudaStreamDestroy(ls_stream[i]);
    delete [] ls_stream;

    return 0;
}

在这段代码中，我们有五个调用，与之前的代码相同，但这里我们将使用五个不同的流。为此，我们建立了一个cudaStream_t数组，并为每个流创建了流。你对这个改变有什么期待？打印输出将与之前的版本相同。运行以下命令来编译这段代码：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_mutli_stream ./2_cuda_multi_stream.cu

然而，这并不保证它们具有相同的操作。正如我们在开始时讨论的，这段代码展示了多个流的并发性，如下面的截图所示：

正如你在截图底部所看到的，五个独立的流同时执行相同的内核函数，并且它们的操作相互重叠。由此，我们可以得出流的两个特点，如下所示：

1. 内核执行与主机是异步的。

2. 不同流中的 CUDA 操作是彼此独立的。

利用流的并发性，我们可以通过重叠独立操作来获得额外的优化机会。

流级别的同步

CUDA 流提供了流级别的同步，使用cudaStreamSynchronize()函数。使用这个函数会强制主机等待直到某个流的操作结束。这为我们迄今为止使用的cudaDeviceSynchronize()函数提供了重要的优化。

我们将在接下来的部分讨论如何利用这一特性，但让我们在这里讨论它的基本操作。前面的例子展示了在循环中没有同步的并发操作。然而，我们可以通过使用cudaStreamSynchronize()函数来阻止主机执行下一个内核执行。下面的代码展示了在内核执行结束时使用流同步的示例：

// execute kernels with the CUDA stream each
for (int i = 0; i < n_stream; i++) {
   foo_kernel<<< 1, 1, 0, ls_stream[i] >>>(i);
   cudaStreamSynchronize(ls_stream[i]);
}

我们可以很容易地预测，由于同步，内核操作的并发性将消失。为了确认这一点，让我们对此进行分析，看看这对内核执行的影响：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_mutli_stream_with_sync ./3_cuda_multi_stream_with_sync.cu

下面的截图显示了结果：

正如你所看到的，所有的内核执行没有重叠点，尽管它们是用不同的流执行的。利用这一特性，我们可以让主机等待特定流操作的开始和结果。

使用默认流

为了让多个流同时运行，我们应该使用我们显式创建的流，因为所有流操作都与默认流同步。下面的截图显示了默认流的同步操作效果：

我们可以通过修改我们的多流内核调用操作来实现这一点，就像这样：

for (int i = 0; i < n_stream; i++)
    if (i == 3)
        foo_kernel<<< 1, 1, 0, 0 >>>(i);
    else
        foo_kernel<<< 1, 1, 0, ls_stream[i] >>>(i);

运行以下命令编译代码：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_multi_stream_with_default ./4_cuda_multi_stream_with_default.cu

因此，我们可以看到最后一个操作无法与前面的内核执行重叠，而是必须等到第四个内核执行完成后才能进行。

GPU 执行的流水线

多个流的主要好处之一是将数据传输与内核执行重叠。通过重叠内核操作和数据传输，我们可以隐藏数据传输开销并提高整体性能。

GPU 流水线的概念

当我们执行内核函数时，我们需要将数据从主机传输到 GPU，然后将结果从 GPU 传输回主机。下面的图表显示了在主机和内核执行之间传输数据的迭代操作的示例：

然而，内核执行基本上是异步的，主机和 GPU 可以同时运行。如果主机和 GPU 之间的数据传输具有相同的特性，我们就能够重叠它们的执行，就像我们在前面的部分中看到的那样。下面的图表显示了当数据传输可以像正常的内核操作一样执行，并与流一起处理时的操作：

在这个图表中，我们可以看到主机和设备之间的数据传输可以与内核执行重叠。然后，这种重叠操作的好处是减少应用程序的执行时间。通过比较两张图片的长度，您将能够确认哪个操作的吞吐量更高。

关于 CUDA 流，所有 CUDA 操作——数据传输和内核执行——在同一个流中是顺序的。然而，它们可以与不同的流同时操作。以下图表显示了多个流的重叠数据传输和内核操作：

为了实现这样的流水线操作，CUDA 有三个先决条件：

1. 主机内存应该分配为固定内存——CUDA 提供了cudaMallocHost()和cudaFreeHost()函数来实现这一目的。

2. 在主机和 GPU 之间传输数据而不阻塞主机——CUDA 提供了cudaMemcpyAsync()函数来实现这一目的。

3. 管理每个操作以及不同的 CUDA 流，以实现并发操作。

现在，让我们编写一个简单的应用程序来对工作负载进行流水线处理。

构建流水线执行

以下代码显示了异步数据传输的片段以及在执行结束时 CUDA 流的同步：

cudaStream_t stream;
float *h_ptr, *d_ptr;    size_t byte_size = sizeof(float) * BUF_SIZE;

cudaStreamCreate(&stream);               // create CUDA stream
cudaMallocHost(h_ptr, byte_size);        // allocates pinned memory
cudaMalloc((void**)&d_ptr, byte_size);   // allocates a global memory

// transfer the data from host to the device asynchronously
cudaMemcpyAsync(d_ptr, h_ptr, byte_size, cudaMemcpyHostToDevice, stream);

... { kernel execution } ...

// transfer the data from the device to host asynchronously
cudaMemcpyAsync(h_ptr, d_ptr, byte_size, cudaMemcpyDeviceToHost, stream);
cudaStreamSynchronize(stream);

// terminates allocated resources
cudaStreamDestroy(stream);
cudaFree(d_ptr);
cudaFreeHost(h_ptr);

这段代码展示了如何分配固定内存，并使用用户创建的流传输数据。通过合并这个例子和多个 CUDA 流操作，我们可以实现流水线 CUDA 操作。

现在，让我们构建一个应用程序，其中包含数据传输和内核执行的流水线操作。在这个应用程序中，我们将使用一个将两个向量相加的内核函数，通过切片流的数量，并输出其结果。然而，内核的实现在主机代码级别不需要任何更改。但是，我们将迭代加法操作 500 次以延长内核执行时间。因此，实现的内核代码如下：

__global__ void
vecAdd_kernel(float *c, const float* a, const float* b)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    for (int i = 0; i < 500; i++)
        c[idx] = a[idx] + b[idx];
}

为了处理每个流的操作，我们将创建一个管理 CUDA 流和 CUDA 操作的类。这个类将允许我们管理 CUDA 流以及索引。以下代码显示了该类的基本架构：

class Operator
{
private:
    int index;

public:
    Operator() {
        cudaStreamCreate(&stream);    // create a CUDA stream
    }

    ~Operator() {
        cudaStreamDestroy(stream);    // terminate the CUDA stream
    }

    cudaStream_t stream;
    void set_index(int idx) { index = idx; }
    void async_operation(float *h_c, const float *h_a, 
                         const float *h_b,
                         float *d_c, float *d_a, float *d_b,
                         const int size, const int bufsize);

}; // Operator

现在，让我们编写一些顺序 GPU 执行代码，这些代码在前一节中已经使用过，但作为Operator类的成员函数，如下所示：

void Operator::async_operation(float *h_c, const float *h_a, 
                          const float *h_b,
                          float *d_c, float *d_a, float *d_b,
                          const int size, const int bufsize)
{
    // start timer
    sdkStartTimer(&_p_timer);

    // copy host -> device
    cudaMemcpyAsync(d_a, h_a, bufsize, 
                    cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_b, h_b, bufsize, 
                    cudaMemcpyHostToDevice, stream);

    // launch cuda kernel
    dim3 dimBlock(256);
    dim3 dimGrid(size / dimBlock.x);
    vecAdd_kernel<<< dimGrid, dimBlock, 0, 
                     stream >>>(d_c, d_a, d_b);

    // copy device -> host
    cudaMemcpyAsync(h_c, d_c, bufsize, 
                    cudaMemcpyDeviceToHost, stream);

    printf("Launched GPU task %d\n", index);
}

这个函数的操作与我们之前使用的基本 CUDA 主机编程模式没有什么不同，只是我们使用了给定的_stream应用了cudaMemcpyAsync()。然后，我们编写main()来处理多个操作符实例和页锁定内存：

int main(int argc, char* argv[])
{
    float *h_a, *h_b, *h_c;
    float *d_a, *d_b, *d_c;
    int size = 1 << 24;
    int bufsize = size * sizeof(float);
    int num_operator = 4;

    if (argc != 1)
        num_operator = atoi(argv[1]);

现在，我们将使用cudaMallocHost()来分配主机内存，以获得固定内存，并对其进行初始化：

    cudaMallocHost((void**)&h_a, bufsize);
    cudaMallocHost((void**)&h_b, bufsize);
    cudaMallocHost((void**)&h_c, bufsize);

    srand(2019);
    init_buffer(h_a, size);
    init_buffer(h_b, size);
    init_buffer(h_c, size);

而且，我们将拥有相同大小的设备内存：

    cudaMalloc((void**)&d_a, bufsize);
    cudaMalloc((void**)&d_b, bufsize);
    cudaMalloc((void**)&d_c, bufsize);

现在，我们将使用我们使用的类创建一个 CUDA 操作符列表：

    Operator *ls_operator = new Operator[num_operator];

我们准备执行流水线操作。在开始执行之前，让我们放一个秒表来查看整体执行时间，并查看重叠数据传输的好处，如下所示：

    StopWatchInterface *timer;
    sdkCreateTimer(&timer);
    sdkStartTimer(&timer);

让我们使用循环执行每个操作符，并且每个操作符将根据其顺序访问主机和设备内存。我们还将测量循环的执行时间：

    for (int i = 0; i < num_operator; i++) {
        int offset = i * size / num_operator;
        ls_operator[i].set_index(i);
        ls_operator[i].async_operation(&h_c[offset], 
                                       &h_a[offset], &h_b[offset],
                                       &d_c[offset], 
                                       &d_a[offset], &d_b[offset],
                                       size / num_operator, 
                                       bufsize / num_operator);
    }

    cudaDeviceSynchronize();
    sdkStopTimer(&timer);

最后，我们将比较一个样本的结果，并打印出整体测量性能：

    // prints out the result
    int print_idx = 256;
    printf("compared a sample result...\n");
    printf("host: %.6f, device: %.6f\n", h_a[print_idx] + 
           h_b[print_idx], h_c[print_idx]);

    // prints out the performance
    float elapsed_time_msed = sdkGetTimerValue(&timer);
    float bandwidth = 3 * bufsize * sizeof(float) / 
                      elapsed_time_msed / 1e6;
    printf("Time= %.3f msec, bandwidth= %f GB/s\n", 
           elapsed_time_msed, bandwidth);

终止句柄和内存，如下所示：

    sdkDeleteTimer(&timer);
    delete [] ls_operator;
    cudaFree(d_a);    cudaFree(d_b);    cudaFree(d_c);
    cudaFreeHost(h_a);cudaFreeHost(h_b);cudaFreeHost(h_c);

要执行代码，让我们重用前面的主机初始化函数和 GPU 内核函数。我们暂时不需要修改这些函数。使用以下命令编译代码：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_pipelining ./cuda_pipelining.cu

您必须使用 GPU 的计算能力版本号作为gencode选项。编译的输出如下：

Launched GPU task 0
Launched GPU task 1
Launched GPU task 2
Launched GPU task 3
compared a sample result...
host: 1.523750, device: 1.523750
Time= 29.508 msec, bandwidth= 27.291121 GB/s

正如我们所看到的，GPU 任务是按照内核执行的顺序以及流的顺序执行的。

现在，让我们来回顾一下应用程序在内部是如何运行的。默认情况下，示例代码将主机数据切片为四个部分，并同时执行四个 CUDA 流。我们可以看到每个核函数的输出以及流的执行情况。要查看重叠操作，您需要使用以下命令对执行进行分析：

$ nvprof -o overlapping_exec.nvvp ./overlapping_exec

以下截图显示了通过重叠数据传输和核函数执行来操作四个 CUDA 流：

核函数执行和数据传输之间的重叠

因此，GPU 可以忙碌直到最后一个核函数执行完成，并且我们可以隐藏大部分数据传输。这不仅增强了 GPU 的利用率，还减少了总应用程序执行时间。

在核函数执行之间，我们可以发现它们虽然属于不同的 CUDA 流，但没有争用。这是因为 GPU 调度器知道执行请求，并优先服务第一个。然而，当当前任务完成时，流多处理器可以为另一个 CUDA 流中的下一个核函数提供服务，因为它们仍然保持占用。

在多个 CUDA 流操作结束时，我们需要同步主机和 GPU，以确认 GPU 上的所有 CUDA 操作都已完成。为此，我们在循环结束后立即使用了cudaDeviceSynchronize()。此函数可以在调用点同步所选的所有 GPU 操作。

对于同步任务，我们可以用以下代码替换cudaDeviceSynchronize()函数。为此，我们还必须将私有成员_stream更改为公共成员：

for (int i = 0; i < num_operator; i++) {
    cudaStreamSynchronize(ls_operator[i]._stream);
}

当我们需要在每个流完成后从单个主机线程提供特定操作时，可以使用这个。但是，这不是一个好的操作设计，因为后续操作无法避免与其他流同步。

在循环中使用cudaStreamSynchronize()怎么样？在这种情况下，我们无法执行之前的重叠操作。以下截图显示了这种情况：

这是因为cudaStreamSynchronize()将同步每次迭代，应用程序将按顺序执行所有 CUDA 执行。在这种情况下，执行时间为 41.521 毫秒，比重叠执行时间慢了约 40%。

CUDA 回调函数

CUDA 回调函数是可调用的主机函数，由 GPU 执行上下文执行。使用此函数，程序员可以指定在 GPU 操作之后执行主机所需的主机操作。

CUDA 回调函数具有一个名为CUDART_CB的特殊数据类型，因此应该使用这种类型进行定义。使用此类型，程序员可以指定哪个 CUDA 流启动此函数，传递 GPU 错误状态，并提供用户数据。

要注册回调函数，CUDA 提供了cudaStreamAddCallback()。该函数接受 CUDA 流、CUDA 回调函数及其参数，以便从指定的 CUDA 流中调用指定的 CUDA 回调函数并获取用户数据。该函数有四个输入参数，但最后一个是保留的。因此，我们不使用该参数，它保持为0。

现在，让我们改进我们的代码，使用回调函数并输出单个流的性能。如果要分开之前的工作和这个工作，可以复制源代码。

首先，将这些函数声明放入Operator类的private区域：

StopWatchInterface *_p_timer;
static void CUDART_CB Callback(cudaStream_t stream, cudaError_t status, void* userData);
void print_time();

Callback()函数将在每个流的操作完成后被调用，print_time()函数将使用主机端计时器_p_timer报告估计的性能。函数的实现如下：

void Operator::CUDART_CB Callback(cudaStream_t stream, cudaError_t status, void* userData) {
    Operator* this_ = (Operator*) userData;
    this_->print_time();
}

void Operator::print_time() {
    sdkStopTimer(&p_timer);    // end timer
    float elapsed_time_msed = sdkGetTimerValue(&p_timer);
    printf("stream %2d - elapsed %.3f ms \n", index, 
           elapsed_time_msed);
}

为了进行正确的计时操作，我们需要在Operator类的构造函数中进行计时器初始化，并在类的终结器中进行计时器销毁。此外，我们必须在Operator::async_operation()函数的开头启动计时器。然后，在函数的末尾插入以下代码块。这允许 CUDA 流在完成先前的 CUDA 操作时调用主机端函数：

// register callback function
cudaStreamAddCallback(stream, Operator::Callback, this, 0);

现在，让我们编译并查看执行结果。您必须使用您的 GPU 的计算能力版本号作为gencode选项：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_callback ./cuda_callback.cu

这是我们更新的执行结果：

stream 0 - elapsed 11.136 ms
stream 1 - elapsed 16.998 ms
stream 2 - elapsed 23.283 ms
stream 3 - elapsed 29.487 ms
compared a sample result...
host: 1.523750, device: 1.523750
Time= 29.771 msec, bandwidth= 27.050028 GB/s

在这里，我们可以看到估计的执行时间以及 CUDA 流。回调函数估计其序列的执行时间。由于与其他流重叠并延迟后续 CUDA 流，我们可以看到后续 CUDA 流的执行时间延长。我们可以通过与分析结果匹配来确认这些经过的时间，如下所示：

尽管它们的测量经过时间随着流的执行而延长，但流之间的差值是固定的，我们可以从分析输出中看到这些操作。

因此，我们可以得出结论，我们可以编写主机代码，以便在每个单独的 CUDA 流操作完成后立即执行。这比从主线程同步每个流更加先进。

具有优先级的 CUDA 流

默认情况下，所有 CUDA 流具有相同的优先级，因此它们可以按正确的顺序执行其操作。此外，CUDA 流还可以具有优先级，并且可以被优先级更高的流取代。有了这个特性，我们可以有满足时间关键要求的 GPU 操作。

CUDA 中的优先级

要使用具有优先级的流，我们首先需要从 GPU 获取可用的优先级。我们可以使用cudaDeviceGetStreamPriorityRange()函数来获取这些值。它的输出是两个数值，即最低和最高的优先级值。然后，我们可以使用cudaStreamCreaetWithPriority()函数创建一个优先级流，如下所示：

cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority)

我们应该提供两个额外的参数。第一个确定了创建的流与默认流的行为。我们可以使用cudaStreamDefault使新流与默认流同步，就像普通流一样。另一方面，我们可以使用cudaStreamNonBlocking使其与默认流并行操作。最后，我们可以在优先级范围内设置流的优先级。在 CUDA 编程中，最低值具有最高优先级。

此外，我们可以使用以下代码确认 GPU 是否支持这一点。但是，我们不必太担心这一点，因为自 CUDA 计算能力 3.5 以来，优先级流一直可用：

cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
if (prop.streamPrioritiesSupported == 0) { ... }

如果设备属性值为0，我们应该停止应用程序，因为 GPU 不支持流优先级。

具有优先级的流执行

现在，我们将重用之前带有回调的多流应用程序。在这段代码中，我们可以看到流可以按顺序操作，我们将看到如何使用优先级更改这个顺序。我们将从Operator类派生一个类，并且它将处理流的优先级。因此，我们将把成员变量流的保护级别从私有成员更改为受保护的成员。构造函数可以选择性地创建流，因为这可以由派生类完成。更改如下代码所示：

... { middle of the class Operator } ...
protected:
    cudaStream_t stream = nullptr;

public:
    Operator(bool create_stream = true) {
        if (create_stream)
            cudaStreamCreate(&stream);
        sdkCreateTimer(&p_timer);
    }
... { middle of the class Operator } ...

派生类Operator_with_priority将具有一个函数，可以根据给定的优先级手动创建一个 CUDA 流。该类的配置如下：

class Operator_with_priority: public Operator {
public:
    Operator_with_priority() : Operator(false) {}

    void set_priority(int priority) {
        cudaStreamCreateWithPriority(&stream, 
            cudaStreamNonBlocking, priority);
    }
};

当我们使用类处理每个流的操作时，我们将更新main()中的ls_operator创建代码，以使用我们之前编写的Operator_with_priority类，如下所示：

Operator_with_priority *ls_operator = new Operator_with_priority[num_operator];

当我们更新类时，这个类在我们请求它之前不会创建流。正如我们之前讨论的，我们需要使用以下代码获取 GPU 可用优先级范围：

// Get priority range
int priority_low, priority_high;
cudaDeviceGetStreamPriorityRange(&priority_low, &priority_high);
printf("Priority Range: low(%d), high(%d)\n", priority_low, priority_high);

然后，让我们创建每个操作以拥有不同的优先级流。为了简化这个任务，我们将让最后一个操作拥有最高的流，并看看 CUDA 流中的抢占是如何工作的。可以使用以下代码来实现这一点：

for (int i = 0; i < num_operator; i++) {
    ls_operator[i].set_index(i);

    // let the latest CUDA stream to have the high priority
    if (i + 1 == num_operator)
        ls_operator[i].set_priority(priority_high);
    else
        ls_operator[i].set_priority(priority_low);
}

之后，我们将执行每个操作，就像之前一样：

for (int i = 0 ; i < num_operator; i++) { 
    int offset = i * size / num_operator;
    ls_operator[i].async_operation(&h_c[offset], 
                                   &h_a[offset], &h_b[offset],
                                   &d_c[offset], 
                                   &d_a[offset], &d_b[offset],
                                   size / num_operator, 
                                   bufsize / num_operator);
}

为了获得正确的输出，让我们使用cudaDeviceSynchronize()函数同步主机和 GPU。最后，我们可以终止 CUDA 流。具有优先级的流可以使用cudaStreamDestroy()函数终止，因此在这个应用程序中我们已经做了必要的事情。

现在，让我们编译代码并查看效果。和往常一样，您需要向编译器提供正确的 GPU 计算能力版本：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o prioritized_cuda_stream ./prioritized_cuda_stream.cu

接下来是应用程序的输出：

Priority Range: low(0), high(-1)
stream 0 - elapsed 11.119 ms
stream 3 - elapsed 19.126 ms
stream 1 - elapsed 23.327 ms
stream 2 - elapsed 29.422 ms
compared a sample result...
host: 1.523750, device: 1.523750
Time= 29.730 msec, bandwidth= 27.087332 GB/s

从输出中，您可以看到操作顺序已经改变。Stream 3 在 Stream 1 和 Stream 2 之前。下面的屏幕截图显示了它是如何改变的：

在这个屏幕截图中，第二个 CUDA 流（在这种情况下是 Stream 19）被优先级最低的 CUDA 流（Stream 21）抢占，以便在 Stream 21 执行完毕后 Stream 19 完成其工作。请注意，数据传输的顺序不会根据这种优先级而改变。

使用 CUDA 事件估计内核执行时间

以前的 GPU 操作时间估计有一个限制，即它无法测量内核执行时间。这是因为我们在主机端使用了计时 API。因此，我们需要与主机和 GPU 同步以测量内核执行时间，考虑到对应用程序性能的开销和影响，这是不切实际的。

这可以通过使用 CUDA 事件来解决。CUDA 事件记录 GPU 端的事件以及 CUDA 流。CUDA 事件可以是基于 GPU 状态的事件，并记录调度时间。使用这个，我们可以触发以下操作或估计内核执行时间。在本节中，我们将讨论如何使用 CUDA 事件测量内核执行时间。

CUDA 事件由cudaEvent_t句柄管理。我们可以使用cudaEventCreate()创建 CUDA 事件句柄，并使用cudaEventDestroy()终止它。要记录事件时间，可以使用cudaEventRecord()。然后，CUDA 事件句柄记录 GPU 的事件时间。这个函数还接受 CUDA 流，这样我们就可以将事件时间枚举到特定的 CUDA 流。在获取内核执行的开始和结束事件之后，可以使用cudaEventElapsedTime()获取经过的时间，单位为毫秒。

现在，让我们讨论如何使用 CUDA 事件来使用这些 API。

使用 CUDA 事件

在本节中，我们将重用第二节中的多流应用程序。然后，我们使用 CUDA 事件枚举每个 GPU 内核的执行时间：

1. 我们将使用一个简单的向量加法内核函数，如下所示：

__global__ void
vecAdd_kernel(float *c, const float* a, const float* b) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    for (int i = 0; i < 500; i++)
        c[idx] = a[idx] + b[idx];
}

这段代码有一个迭代，它延长了内核执行时间。

1. 然后，我们将使用以下片段来测量内核执行时间。为了比较结果，我们将使用主机端的计时器和 CUDA 事件：

... { memory initializations } ...

// initialize the host timer
StopWatchInterface *timer;
sdkCreateTimer(&timer);

cudaEvent_t start, stop;
// create CUDA events
cudaEventCreate(&start);
cudaEventCreate(&stop);

// start to measure the execution time
sdkStartTimer(&timer);
cudaEventRecord(start);

// launch cuda kernel
dim3 dimBlock(256);
dim3 dimGrid(size / dimBlock.x);
vecAdd_kernel<<< dimGrid, dimBlock >>>(d_c, d_a, d_b);

// record the event right after the kernel execution finished
cudaEventRecord(stop);

// Synchronize the device to measure the execution time from the host side
cudaEventSynchronize(stop); // we also can make synchronization based on CUDA event
sdkStopTimer(&timer);

正如您在这段代码中所看到的，我们可以在内核调用之后立即记录 CUDA 事件。然而，计时器需要在 GPU 和主机之间进行同步。为了同步，我们使用cudaEventSynchronize(stop)函数，因为我们也可以使主机线程与事件同步。与此同时，这段代码只涵盖了处理计时资源和内核执行。但是，您还需要初始化所需的内存才能使其工作。

1. 在内核执行之后，让我们编写代码报告每个计时资源的执行时间：

// print out the result
int print_idx = 256;
printf("compared a sample result...\n");
printf("host: %.6f, device: %.6f\n", h_a[print_idx] + h_b[print_idx], h_c[print_idx]);

// print estimated kernel execution time
float elapsed_time_msed = 0.f;
cudaEventElapsedTime(&elapsed_time_msed, start, stop);
printf("CUDA event estimated - elapsed %.3f ms \n", elapsed_time_msed);

1. 现在，我们将通过终止计时资源来完成我们的应用程序，使用以下代码：

// delete timer
sdkDeleteTimer(&timer);

// terminate CUDA events
cudaEventDestroy(start);
cudaEventDestroy(stop);

1. 让我们编译并使用以下命令查看输出：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_event ./cuda_event.cu
compared a sample result...
host: 1.523750, device: 1.523750
CUDA event estimated - elapsed 23.408 ms 
Host measured time= 35.063 msec/s

如您所见，我们可以使用 CUDA 事件来测量内核执行时间。但是，测量的时间在 CUDA 事件和计时器之间存在间隙。我们可以使用 NVIDIA 分析器来验证哪个提供更准确的信息。当我们使用# nvprof ./cuda_event命令时，输出如下：

如您所见，与从主机测量相比，CUDA 事件提供了准确的结果。

使用 CUDA 事件的另一个好处是，我们可以使用多个 CUDA 流同时测量多个内核执行时间。让我们实现一个示例应用程序并查看其操作。

多流估计

cudaEventRecord()函数对主机是异步的。换句话说，没有同步来测量内核执行时间到示例代码。为了使事件和主机同步，我们需要使用cudaEventSynchronize()。例如，当我们在cudaEventRecord(stop)之后立即放置这个函数时，可以在设备到主机的异步数据传输之前放置内核函数打印，通过同步效果来实现。

在多个 CUDA 流应用程序中测量内核执行时间也是有用的：

1. 让我们将这应用到04_stream_priority示例代码中的多个 CUDA 流重叠的代码中。使用以下代码更新代码：

class Operator
{
private:
    int _index;
    cudaStream_t stream;
    StopWatchInterface *p_timer;
    cudaEvent_t start, stop;

public:
    Operator() {
        cudaStreamCreate(&stream);

 // create cuda event
 cudaEventCreate(&start);
 cudaEventCreate(&stop);
    }

    ~Operator() {
        cudaStreamDestroy(stream);

 // destroy cuda event
 cudaEventDestroy(start);
 cudaEventDestroy(stop);
    }

    void set_index(int idx) { index = idx; }
    void async_operation(float *h_c, const float *h_a, 
                          const float *h_b,
                          float *d_c, float *d_a, float *d_b,
                          const int size, const int bufsize);
 void print_kernel_time();

}; // Operator

1. 然后，我们将定义此时包含的print_time()函数，如下所示：

void Operator::print_time() {
    float milliseconds = 0;
    cudaEventElapsedTime(&milliseconds, start, stop);
    printf("Stream %d time: %.4f ms\n", index, milliseconds);
}

1. 现在，在Operator::async_operation()的开头和结尾插入cudaEventRecord()函数调用，如下所示：

void Operator::async_operation( ... )
{
    // start timer
    sdkStartTimer(&p_timer);

    // copy host -> device
    cudaMemcpyAsync(d_a, h_a, bufsize, 
                    cudaMemcpyHostToDevice, stream);
    cudaMemcpyAsync(d_b, h_b, bufsize, 
                    cudaMemcpyHostToDevice, stream);

    // record the event before the kernel execution
 cudaEventRecord(start, stream);

    // launch cuda kernel
    dim3 dimBlock(256);
    dim3 dimGrid(size / dimBlock.x);
    vecAdd_kernel<<< dimGrid, dimBlock, 0, 
                     stream >>>(d_c, d_a, d_b);

    // record the event right after the kernel execution finished
 cudaEventRecord(stop, stream);

    // copy device -> host
    cudaMemcpyAsync(h_c, d_c, bufsize, 
                    cudaMemcpyDeviceToHost, stream);

    // what happen if we include CUDA event synchronize?
    // QUIZ: cudaEventSynchronize(stop);

    // register callback function
    cudaStreamAddCallback(stream, Operator::Callback, this, 0);
}

对于这个函数，在函数的末尾放置同步是一个挑战。在完成本节后尝试这样做。这将影响应用程序的行为。建议尝试自己解释输出，然后使用分析器进行确认。

现在，让我们编译并查看执行时间报告，如下；它显示与先前执行类似的性能：

$ nvcc -m64 -run -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o cuda_event_with_streams ./cuda_event_with_streams.cu
Priority Range: low(0), high(-1)
stream 0 - elapsed 11.348 ms 
stream 3 - elapsed 19.435 ms 
stream 1 - elapsed 22.707 ms 
stream 2 - elapsed 35.768 ms 
kernel in stream 0 - elapsed 6.052 ms 
kernel in stream 1 - elapsed 14.820 ms 
kernel in stream 2 - elapsed 17.461 ms 
kernel in stream 3 - elapsed 6.190 ms 
compared a sample result...
host: 1.523750, device: 1.523750
Time= 35.993 msec, bandwidth= 22.373972 GB/s

在这个输出中，我们还可以看到每个内核的执行时间，这要归功于 CUDA 事件。从这个结果中，我们可以看到内核执行时间延长了，就像我们在上一节中看到的那样。

如果您想了解更多关于 CUDA 事件特性的信息，请查看 NVIDIA 的 CUDA 事件文档：docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__EVENT.html。

现在，我们将介绍管理 CUDA 网格的其他一些方面。第一项是动态并行性，它使 GPU 内核函数能够进行内核调用。

CUDA 动态并行性

CUDA 动态并行性（CDP）是一种设备运行时功能，它允许从设备函数进行嵌套调用。这些嵌套调用允许子网格具有不同的并行性。当问题需要不同的块大小时，此功能非常有用。

理解动态并行性

与主机的普通内核调用一样，GPU 内核调用也可以进行内核调用。以下示例代码显示了它的工作原理：

__global__ void child_kernel(int *data) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    atomicAdd(&data[idx], seed);
}

__global__ void parent_kernel(int *data)
{
 if (threadIdx.x == 0) {
        int child_size = BUF_SIZE/gridDim.x;
        child_kernel<<< child_size/BLOCKDIM, BLOCKDIM >>>
                        (&data[child_size*blockIdx.x], blockIdx.x+1);
    }
    // synchronization for other parent's kernel output
    cudaDeviceSynchronize();
}

如您在这些函数中所见，我们需要确保哪个 CUDA 线程进行内核调用以控制网格创建的数量。要了解更多信息，让我们使用这个实现第一个应用程序。

动态并行性的使用

我们的动态并行性代码将创建一个父网格，该父网格将创建一些子网格：

1. 首先，我们将使用以下代码编写parent_kernel()函数和child_kernel()函数：

#define BUF_SIZE (1 << 10)
#define BLOCKDIM 256

__global__ void child_kernel(int *data)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    atomicAdd(&data[idx], 1);
}

__global__ void parent_kernel(int *data)
{
    if (blockIdx.x * blockDim.x + threadIdx.x == 0)
    {
        int child_size = BUF_SIZE/gridDim.x;
        child_kernel<<< child_size/BLOCKDIM, BLOCKDIM >>> \
                        (&data[child_size*blockIdx.x], 
                         blockIdx.x+1);
    }
    // synchronization for other parent's kernel output
    cudaDeviceSynchronize();
}

如您在这段代码中所见，父内核函数创建子内核网格作为块的数量。然后，子网格递增指定的内存1来标记它们的操作。内核执行后，父内核使用cudaDeviceSynchronize()函数等待所有子网格完成其工作。在进行同步时，我们应确定同步的范围。如果我们需要在块级别进行同步，我们应选择__synchthread()。

1. 使用以下代码编写main()函数：

#define BUF_SIZE (1 << 10)
#define BLOCKDIM 256
int main()
{
    int *data;
    int num_child = 4;

    cudaMallocManaged((void**)&data, BUF_SIZE * sizeof(int));
    cudaMemset(data, 0, BUF_SIZE * sizeof(int));

    parent_kernel<<<num_child, 1>>>(data);
    cudaDeviceSynchronize();

    // Count elements value
    int counter = 0;
    for (int i = 0; i < BUF_SIZE; i++)
        counter += data[i];

    // getting answer
    int counter_h = 0;
    for (int i = 0; i < num_child; i++)
        counter_h += (i+1);
    counter_h *= BUF_SIZE / num_child;

    if (counter_h == counter)
        printf("Correct!!\n");
    else
        printf("Error!! Obtained %d. It should be %d\n", 
               counter, counter_h);

    cudaFree(data);
    return 0;
}

正如前面讨论的，我们将创建子网格以及块的数量。因此，我们将使用网格大小为4来执行父内核函数，而块大小为1。

1. 要编译 CDP 应用程序，我们应该为nvcc编译器提供-rdc=true选项。因此，编译源代码的命令如下：

$ nvcc -run -rdc=true -lcudadevrt -gencode arch=compute_70,code=sm_70 -o host_callback host_callback.cu -I/usr/local/cuda/samples/common/inc 

1. 让我们对这个应用程序进行分析，以了解其操作。以下截图显示了这个嵌套调用的工作原理：

如我们在这个屏幕截图中所见，父内核创建了一个子网格，我们可以在左侧面板的右角标中看到它们的关系。然后，父网格（parent_kernel）等待其执行，直到子网格完成其工作。CUDA 目前不支持 SM70（Volta 架构）的 CDT 分析，因此我使用 Tesla P40 来获得这个输出。

递归

动态并行性的一个好处是我们可以创建递归。以下代码显示了一个递归内核函数的示例：

__global__ void recursive_kernel(int *data, int size, int depth) {
  int x_0 = blockIdx.x * size;

  if (depth > 0) {
    __syncthreads();
 if (threadIdx.x == 0) {
        int dimGrid = size / dimBlock;
        recursive_kernel<<<dimGrid, 
              dimBlock>>>(&data[x_0], size/dimGrid, depth-1);
        cudaDeviceSynchronize();
      }
      __syncthreads();
   }
}

如您所见，与以前的动态并行内核函数相比，没有太大的区别。但是，我们应该谨慎使用这个功能，考虑到资源使用和限制。一般来说，动态并行内核可以保守地保留高达 150MB 的设备内存来跟踪待处理的网格启动和通过在子网格启动上进行同步来同步父网格的状态。此外，同步必须在多个级别上小心进行，而嵌套内核启动的深度限制为 24 级。最后，控制嵌套内核启动的运行时可能会影响整体性能。

如果您需要了解动态并行性的限制和限制，请参阅以下编程指南：docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#implementation-restrictions-and-limitations。

我们将在第七章中介绍其在 CUDA 中快速排序实现中的应用，即CUDA 中的并行编程模式。要了解更多关于动态并行性的信息，请参阅以下文档：

• devblogs.nvidia.com/cuda-dynami…

• on-demand.gputechconf.com/gtc/2012/presentations/S0338-New-Features-in-the-CUDA-Programming-Model.pdf

网格级别的协作组

如 第三章 中所讨论的，CUDA 提供了协作组。协作组可以根据其分组目标进行分类：warp 级别、块级别和网格级别的组。本文介绍了网格级别的协作组，并探讨了协作组如何处理 CUDA 网格。

协作组最显著的好处是对目标并行对象的显式同步。使用协作组，程序员可以设计他们的应用程序来显式同步 CUDA 并行对象、线程块或网格。使用第三章中介绍的块级协作组，CUDA 线程编程，我们可以通过指定需要同步的 CUDA 线程或块来编写更易读的代码。

理解网格级协作组

自 9.0 版本以来，CUDA 提供了另一级协作组，与网格一起工作。具体来说，有两个网格级协作组：grid_group和multi_grid_group。使用这些组，程序员可以描述网格在单个 GPU 或多个 GPU 上的操作同步。

在这个示例中，我们将探索grid_group的功能，它可以同步网格与减少问题，就像第三章中所提到的，CUDA 线程编程，关于基于块级减少的先前减少设计。每个线程块产生自己的减少结果，并将它们存储到全局内存中。然后，另一个块级减少内核启动，直到我们获得单个减少值。这是因为完成内核操作可以保证下一个减少内核从多个线程块中读取减少值。其设计由左侧的图表描述：

！

另一方面，网格级同步使另一种内部同步块式减少结果的内核设计成为可能，以便主机只需调用一次内核即可获得减少结果。在协作组中，grid_group.sync()提供了这样的功能，因此我们可以编写减少内核而无需内核级迭代。

要使用grid_group.sync()函数，我们需要使用cudaLaunchCooperativeKernel()函数调用内核函数。其接口设计如下：

__host__ cudaError_t cudaLaunchCooperativeKernel
    ( const T* func, dim3 gridDim, dim3 blockDim, 
      void** args, size_t sharedMem = 0, cudaStream_t stream = 0 )

因此，它的使用方式与cudaLaunchKernel()函数相同，该函数启动内核函数。

为了使grid_group中的所有线程块同步，网格中活动线程块的总数不应超过内核函数和设备的最大活动块数。GPU 上的最大活动块大小是每个 SM 的最大活动块数和流处理器的数量的乘积。违反此规则可能导致死锁或未定义行为。我们可以使用cudaOccupancyMaxActiveBlocksPerMultiprocessor()函数来获取每个 SM 内核函数的最大活动线程块数，通过传递内核函数和块大小信息。

使用grid_group的用法

现在，让我们将grid_group应用于并行减少问题，并看看 GPU 编程如何改变：

1. 我们将重用之前并行减少代码中的主机代码，即03_cuda_thread_programming/07_cooperative_groups。换句话说，我们将通过对主机代码进行小的更改来改变 GPU 的操作。您还可以使用07_grid_level_cg目录中的代码。

2. 现在，让我们编写一些块级减少代码。当我们有网格级协作组时，所有线程块必须是活动的。换句话说，我们不能执行多个线程块，而 GPU 能够执行的活动块。因此，这个减少将首先累积输入数据，以覆盖所有数据，使用有限数量的线程块。然后，它将在块级进行并行减少，就像我们在第三章中所介绍的那样，CUDA 线程编程。

以下代码显示了它的实现：

__device__ void
block_reduction(float *out, float *in, float *s_data, int active_size, int size, 
          const cg::grid_group &grid, const cg::thread_block &block)
{
  int tid = block.thread_rank();

  // Stride over grid and add the values to a shared memory buffer
  s_data[tid] = 0.f;
  for (int i = grid.thread_rank(); i < size; i += active_size)
    s_data[tid] += in[i];

  block.sync();

  for (unsigned int stride = blockDim.x / 2; 
       stride > 0; stride >>= 1) {
    if (tid < stride)
      s_data[tid] += s_data[tid + stride];
    block.sync();
  }

  if (block.thread_rank() == 0)
    out[block.group_index().x] = s_data[0];
}

1. 然后，让我们编写一个内核函数，考虑活动块数和grid_group执行块级减少。在这个函数中，我们将调用块级减少代码，并在网格级别进行同步。然后，我们将从输出中执行并行减少，就像我们在第三章 CUDA 线程编程中所介绍的那样。以下代码显示了其实现：

__global__ void
reduction_kernel(float *g_out, float *g_in, unsigned int size)
{
  cg::thread_block block = cg::this_thread_block();
  cg::grid_group grid = cg::this_grid();
  extern __shared__ float s_data[];

  // do reduction for multiple blocks
  block_reduction(g_out, g_in, s_data, grid.size(), 
                  size, grid, block);

  grid.sync();

  // do reduction with single block
  if (block.group_index().x == 0)
    block_reduction(g_out, g_out, s_data, block.size(), gridDim.x, grid, block);
}

1. 最后，我们将实现调用具有可用活动线程块维度的内核函数的主机代码。为此，此函数使用cudaoccupancyMaxActiveBlocksPerMultiprocessor()函数。此外，网格级合作组要求我们通过cudaLaunchCooperativeKernel()函数调用内核函数。您可以在这里看到实现：

int reduction_grid_sync(float *g_outPtr, float *g_inPtr, int size, int n_threads)
{ 
  int num_blocks_per_sm;
  cudaDeviceProp deviceProp;

  // Calculate the device occupancy to know 
  // how many blocks can be run concurrently
  cudaGetDeviceProperties(&deviceProp, 0);
  cudaOccupancyMaxActiveBlocksPerMultiprocessor(&num_blocks_per_sm, 
      reduction_kernel, n_threads, n_threads*sizeof(float));
  int num_sms = deviceProp.multiProcessorCount;
  int n_blocks = min(num_blocks_per_sm * num_sms, 
                     (size + n_threads - 1) / n_threads);

  void *params[3];
  params[0] = (void*)&g_outPtr;
  params[1] = (void*)&g_inPtr;
  params[2] = (void*)&size;
  cudaLaunchCooperativeKernel((void*)reduction_kernel, 
                              n_blocks, n_threads, params, 
                              n_threads * sizeof(float), NULL);

  return n_blocks;
}

1. 现在，请确保可以从reduction.cpp文件中调用主机函数。

2. 然后，让我们编译代码并查看其操作。以下 shell 命令编译代码并执行应用程序。计算能力应该等于或大于70：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -rdc=true -o reduction ./reduction.cpp ./reduction_kernel.cu
Time= 0.474 msec, bandwidth= 141.541077 GB/s
host: 0.996007, device 0.996007

输出性能远远落后于我们在第三章 CUDA 线程编程的最终结果。由于block_reduction()函数在开始时使用了高内存吞吐量，因此它是高度内存绑定的：

主要影响因素是我们只能使用活动线程块。因此，我们无法隐藏内存访问时间。实际上，使用grid_group还有其他目的，例如图搜索、遗传算法和粒子模拟，这要求我们保持状态长时间处于活动状态以获得性能。

这种网格级同步可以为性能和可编程性提供更多好处。由于这使得内核可以自行同步，我们可以使内核自行迭代。因此，它对解决图搜索、遗传算法和实际模拟非常有用。要了解有关grid_groups中合作组的更多信息，请参阅提供的文档on-demand.gputechconf.com/gtc/2017/presentation/s7622-Kyrylo-perelygin-robust-and-scalable-cuda.pdf。

使用 OpenMP 的 CUDA 内核调用

为了增加应用程序的并发性，我们可以从主机的并行任务中进行内核调用。例如，OpenMP 提供了多核架构的简单并行性。本教程介绍了 CUDA 如何操作 OpenMP。

OpenMP 和 CUDA 调用

OpenMP 使用分叉-合并模型的并行性来针对多核 CPU。主线程启动并行操作并创建工作线程。主机线程并行运行自己的工作，并在完成工作后加入。

使用 OpenMP，CUDA 内核调用可以与多个线程并行执行。这有助于程序员不必维护单独的内核调用，而是允许它们的内核执行依赖于主机线程的索引。

在本节中，我们将使用以下 OpenMP API：

• omp_set_num_threads()设置将并行工作的工作线程数。

• omp_get_thread_num()返回工作线程的索引，以便每个线程可以识别其任务。

• #pragma omp parallel {} 指定了一个并行区域，将由工作线程覆盖。

现在，让我们编写一些代码，其中 OpenMP 调用 CUDA 内核函数。

CUDA 与 OpenMP 的内核调用

在本节中，我们将实现一个使用 OpenMP 的多流矢量加法应用程序。为此，我们将修改先前的版本并查看差异：

1. 要测试 CUDA 中的 OpenMP，我们将修改03_cuda_callback目录中的代码。我们将修改main()函数的主体，或者您可以使用放置在08_openmp_cuda目录中的提供的示例代码。

2. 现在，让我们包括 OpenMP 头文件并修改代码。要在代码中使用 OpenMP，我们应该使用#include <omp.h>。而且，我们将更新代码，使其使用 OpenMP 来迭代每个流：

// execute each operator collesponding data
omp_set_num_threads(num_operator);
#pragma omp parallel
{
    int i = omp_get_thread_num();
    printf("Launched GPU task %d\n", i);

    int offset = i * size / num_operator;
    ls_operator[i].set_index(i);
    ls_operator[i].async_operation(&h_c[offset], &h_a[offset],   
                                   &h_b[offset],&d_c[offset], 
                                   &d_a[offset], &d_b[offset],
                                   size / num_operator, bufsize 
                                   / num_operator);
}

1. 使用以下命令编译代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -Xcompiler -fopenmp -lgomp -o openmp ./openmp.cu
stream 0 - elapsed 10.734 ms 
stream 2 - elapsed 16.153 ms 
stream 3 - elapsed 21.968 ms 
stream 1 - elapsed 27.668 ms 
compared a sample result...
host: 1.523750, device: 1.523750
Time= 27.836 msec, bandwidth= 28.930389 GB/s

每当您执行此应用程序时，您将看到每个流以无序方式完成其工作。此外，每个流显示不同的时间。这是因为 OpenMP 可以创建多个线程，并且操作是在运行时确定的。

为了了解其运行情况，让我们对应用程序进行分析。以下截图显示了应用程序的分析时间表。由于调度的原因，这可能与您的情况不同：

如您在此截图中所见，您将能够看到数据传输与 Stream 17 相比已经反转。因此，我们可以看到第二个流最终完成了它的工作。

多进程服务

GPU 能够从并发的 CPU 进程中执行内核。但是，默认情况下，它们只以分时方式执行，即使每个内核没有充分利用 GPU 计算资源。为了解决这种不必要的串行化，GPU 提供了多进程服务（MPS）模式。这使得不同的进程能够同时在 GPU 上执行它们的内核，以充分利用 GPU 资源。启用时，nvidia-cuda-mps-control守护进程监视目标 GPU，并使用该 GPU 管理进程内核操作。此功能仅在 Linux 上可用。在这里，我们可以看到多个进程共享同一个 GPU 的 MPS：

正如我们所看到的，每个进程在 GPU 上并行运行一部分（绿色条），而一部分在 CPU 上运行（蓝色条）。理想情况下，您需要蓝色条和绿色条都能获得最佳性能。这可以通过利用所有最新 GPU 支持的 MPS 功能来实现。

请注意，当一个 MPI 进程无法饱和整个 GPU 并且代码的重要部分也在 CPU 上运行时，多个 MPI 进程在同一个 GPU 上运行是有益的。如果一个 MPI 进程利用整个 GPU，即使 CPU 部分（蓝色条）会减少，绿色条的时间也不会减少，因为 GPU 完全被一个 MPI 进程利用。其他 MPI 进程将根据 GPU 架构以分时方式依次访问 GPU。这类似于启动并发内核的情况。如果一个内核利用整个 GPU，那么另一个内核要么等待第一个内核完成，要么进行分时。

这样做的好处是不需要对应用程序进行任何更改即可使用 MPS。MPS 进程作为守护进程运行，如下命令所示：

$nvidia-smi -c EXCLUSIVE_PROCESS 
$nvidia-cuda-mps-control –d

运行此命令后，所有进程都将其命令提交给 MPS 守护进程，该守护进程负责将 CUDA 命令提交给 GPU。对于 GPU，只有一个进程访问 GPU（MPS 守护进程），因此多个进程可以同时运行来自多个进程的多个内核。这可以帮助将一个进程的内存复制与其他 MPI 进程的内核执行重叠。

消息传递接口简介

消息传递接口（MPI）是一种并行计算接口，它能够触发多个进程跨计算单元 - CPU 核心、GPU 和节点。典型的密集多 GPU 系统包含 4-16 个 GPU，而 CPU 核心的数量在 20-40 个之间。在启用 MPI 的代码中，应用程序的某些部分作为不同的 MPI 进程在多个核心上并行运行。每个 MPI 进程都将调用 CUDA。了解将 MPI 进程映射到相应的 GPU 非常重要。最简单的映射是 1:1，即每个 MPI 进程都独占相应的 GPU。此外，我们还可以将多个 MPI 进程理想地映射到单个 GPU 上。

为了将多进程应用场景应用到单个 GPU 上，我们将使用 MPI。要使用 MPI，您需要为您的系统安装 OpenMPI。按照以下步骤在 Linux 上安装 OpenMPI。此操作已在 Ubuntu 18.04 上进行了测试，因此如果您使用其他发行版，可能会有所不同：

$ wget -O /tmp/openmpi-3.0.4.tar.gz https://www.open-mpi.org/software/ompi/v3.0/downloads/openmpi-3.0.4.tar.gz
$ tar xzf /tmp/openmpi-3.0.4.tar.gz -C /tmp
$ cd /tmp/openmpi-3.0.4
$ ./configure --enable-orterun-prefix-by-default --with-cuda=/usr/local/cuda
$ make -j $(nproc) all && sudo make install
$ sudo ldconfig
$ mpirun --version
mpirun (Open MPI) 3.0.4

Report bugs to http://www.open-mpi.org/community/help/

现在，让我们实现一个可以与 MPI 和 CUDA 一起工作的应用程序。

实现一个启用 MPI 的应用程序

要使应用程序与 MPI 一起工作，我们需要在应用程序中放入一些可以理解 MPI 命令的代码：

1. 我们将重用 OpenMP 示例代码，因此将openmp.cu文件复制到08_openmp_cuda目录中。

2. 在代码开头插入mpi头文件include语句：

#include <mpi.h>

1. 在main()函数中创建秒表后立即插入以下代码：

// set num_operator as the number of requested process
int np, rank;
MPI_Init(&argc, &argv);
MPI_Comm_size(MPI_COMM_WORLD, &np);
MPI_Comm_rank(MPI_COMM_WORLD, &rank);

1. 按照第 3 步中提到的代码后，将所需的内存大小切割为进程数，如下所示：

bufsize /= np;
size /= np;

1. 我们需要让每个线程报告它们所属的进程。让我们更新并行执行代码块中的printf()函数，如下所示：

// execute each operator collesponding data
omp_set_num_threads(num_operator);
#pragma omp parallel
{
    int i = omp_get_thread_num();
    int offset = i * size / num_operator;
    printf("Launched GPU task (%d, %d)\n", rank, i);

    ls_operator[i].set_index(i);
    ls_operator[i].async_operation(&h_c[offset], 
                                   &h_a[offset], &h_b[offset],
                                   &d_c[offset], &d_a[offset], 
                                   &d_b[offset],
                                   size / num_operator, 
                                   bufsize / num_operator);
}

1. 在main()的末尾放置MPI_Finalize()函数以关闭 MPI 实例。

2. 使用以下命令编译代码：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -I/usr/local/include/ -Xcompiler -fopenmp -lgomp -lmpi -o simpleMPI ./simpleMPI.cu

您必须使用 GPU 的计算能力版本号来选择gencode选项。

1. 使用以下命令测试编译后的应用程序：

$ ./simpleMPI 2

1. 现在，使用以下命令测试 MPI 执行：

$ mpirun -np 2 ./simpleMPI 2
Number of process: 2
Number of operations: 2
Launched GPU task (1, 0)
Launched GPU task (1, 1)
Number of operations: 2
Launched GPU task (0, 0)
Launched GPU task (0, 1)
stream 0 - elapsed 13.390 ms 
stream 1 - elapsed 25.532 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 25.749 msec, bandwidth= 15.637624 GB/s
stream 0 - elapsed 21.334 ms 
stream 1 - elapsed 26.010 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 26.111 msec, bandwidth= 15.420826 GB/s

启用 MPS

在 GPU 上启用 MPS 需要对 GPU 操作模式进行一些修改。但是，您需要具有比 Kepler 架构更晚的 GPU 架构。

让我们按照以下步骤启用 MPS：

1. 使用以下命令启用 MPS 模式：

$ export CUDA_VISIBLE_DEVICES=0
$ sudo nvidia-smi -i 0 -c 3
$ sudo nvidia-cuda-mps-control -d

或者，您可以使用make enable_mps命令来使用此预定义在Makefile中的配方示例代码。然后，我们可以从nivida-smi输出中看到更新后的计算模式：

1. 现在，使用以下命令测试 MPS 模式下的 MPI 执行：

$ mpirun -np 2 ./simpleMPI 2
Number of process: 2
Number of operations: 2
Launched GPU task (1, 0)
Launched GPU task (1, 1)
stream 0 - elapsed 10.203 ms 
stream 1 - elapsed 15.903 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 16.129 msec, bandwidth= 24.964548 GB/s
Number of operations: 2
Launched GPU task (0, 0)
Launched GPU task (0, 1)
stream 0 - elapsed 10.203 ms 
stream 1 - elapsed 15.877 ms 
compared a sample result...
host: 1.306925, device: 1.306925
Time= 15.997 msec, bandwidth= 25.170544 GB/s

如您所见，与之前的执行相比，每个进程的经过时间都有所减少。

1. 现在，让我们恢复原始模式。要禁用 MPS 模式，请使用以下命令：

$ echo "quit" | sudo nvidia-cuda-mps-control
$ sudo nvidia-smi -i 0 -c 0

或者，您可以使用make disable_mps命令来使用此预定义在Makefile中的配方示例代码。

要了解更多关于 MPS 的信息，请使用以下链接：

• on-demand.gputechconf.com/gtc/2015/presentation/S5584-Priyanka-Sah.pdf

• docs.nvidia.com/deploy/pdf/CUDA_Multi_Process_Service_Overview.pdf

对 MPI 应用程序进行分析并了解 MPS 操作

使用 MPI，多个进程的内核可以同时共享 GPU 资源，从而增强整体 GPU 利用率。没有 MPS，由于时间切片共享和上下文切换开销，GPU 资源被低效地共享。

以下屏幕截图显示了没有 MPS 的多个进程的时间轴配置文件结果：

在此配置文件中，我们可以看到两个 CUDA 上下文共享一个 GPU，并且由于上下文之间的时间共享，内核执行时间延长。

另一方面，MPS 模式管理内核执行请求，因此所有内核执行都会像使用单个进程一样启动。以下屏幕截图显示了 MPS 模式下的内核执行：

如您所见，只有一个 CUDA 流驻留在 GPU 上并控制所有 CUDA 流。此外，所有内核执行时间都得到了稳定，并且使用 MPS 可以减少总的经过时间。总之，使用 MPS 模式有利于多个 GPU 进程的整体性能，并共享 GPU 资源。

nvprof支持将多个 MPI 进程的分析器信息转储到不同的文件中。例如，对于基于 Open MPI 的应用程序，以下命令将在多个文件中转储分析信息，每个文件的名称都基于 MPI 进程的排名：

$ mpirun -np 2 nvprof -f -o simpleMPI.%q{OMPPI_COMM_WORLD_RANK}_2.nvvp ./simpleMPI 2

或者，您可以使用以下命令来执行示例代码：

$ PROCS=2 STREAMS=2 make nvprof

然后，您将为每个进程获得两个nvvp文件。

现在，我们将使用以下步骤使用 NVIDIA Visual Profiler 来查看这些nvvp文件：

1. 打开文件|导入菜单，通过导入nvvp文件创建一个分析会话：

在 Windows 或 Linux 中，快捷键是Ctrl + I，OSX 使用command + I。

1. 然后从列表中选择 Nvprof 后，点击下一步按钮：

1. 从 Nvprof 选项中，选择多个进程，然后单击下一步>：

1. 从导入 Nvprof 数据中，单击浏览...按钮，并选择由nvprof生成的nvvp文件。要对具有多个进程的应用程序进行分析，您需要导入nvvp文件，因为存在多个进程：

1. 单击完成，然后 NVIDIA Visual Profiler 将以时间线视图显示分析结果，如下所示：

请注意，只有同步 MPI 调用将由nvprof进行注释。如果使用异步 MPI API，则需要使用其他 MPI 专用的分析工具。其中一些最著名的工具包括以下内容：

• TAU：TAU 是一种性能分析工具包，目前由俄勒冈大学维护。

• Vampir：这是一种商业可用的工具，对数百个 MPI 进程具有良好的可伸缩性。

• Intel VTune Amplifier：商业工具的另一个选择是 Intel VTune Amplifier。它是目前可用的最好的工具之一，可用于 MPI 应用程序分析。

最新的 CUDA 工具包还允许对 MPI API 进行注释。为此，需要将--annotate-mpi标志传递给nvprof，如以下命令所示：

mpirun -np 2 nvprof --annotate-mpi openmpi -o myMPIApp.%q{OMPI_COMM_WORLD_RANK}.nvprof ./myMPIApplciation

内核执行开销比较

对于迭代并行 GPU 任务，我们有三种内核执行方法：迭代内核调用，具有内部循环，以及使用动态并行性进行递归。最佳操作由算法和应用程序确定。但是，您也可以考虑它们之间的内核执行选项。本示例帮助您比较这些内核执行开销并审查它们的可编程性。

首先，让我们确定我们将测试哪种操作。本示例将使用一个简单的 SAXPY 操作。这有助于我们专注并制作迭代执行代码。此外，随着操作变得更简单，操作控制开销将变得更重。但是，您当然可以尝试任何其他操作。

实现三种内核执行方式

以下步骤涵盖了三种不同迭代操作的性能比较：

1. 创建并导航到10_kernel_execution_overhead目录。

2. 编写simple_saxpy_kernel()函数，代码如下：

__global__ void
simple_saxpy_kernel(float *y, const float* x, const float alpha, const float beta)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    y[idx] = alpha * x[idx] + beta;
}

1. 编写iterative_saxpy_kernel()函数，代码如下：

__global__ void
iterative_saxpy_kernel(float *y, const float* x, 
                       const float alpha, const float beta, 
                       int n_loop)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    for (int i = 0; i < n_loop; i++)
        y[idx] = alpha * x[idx] + beta;
}

1. 编写recursive_saxpy_kernel()函数，代码如下：

__global__ void
recursive_saxpy_kernel(float *y, const float* x, 
                       const float alpha, const float beta, 
                       int depth)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;

    if (depth == 0)
        return;
    else
        y[idx] = alpha * x[idx] + beta;

    if (idx == 0)
        vecAdd_kernel_C<<< gridDim.x, blockDim.x 
                           >>>(y, x, alpha, beta, depth - 1);
}

1. 编写启动这些 CUDA 内核函数的主机代码。首先，我们将对simple_saxpy_kernel()函数进行迭代调用：

for (int i = 0; i < n_loop; i++) {
    simple_saxpy_kernel<<< dimGrid, dimBlock >>>(
                           d_y, d_x, alpha, beta);
}

其次，我们将调用iterative_saxpy_kernel()内核函数，该函数内部有一个迭代循环：

iterative_saxpy_kernel<<< dimGrid, dimBlock >>>(
                          d_y, d_x, alpha, beta, n_loop);

最后，我们将调用recursive_saxpy_kernel()内核函数，该函数以递归方式调用自身：

recursive_saxpy_kernel<<< dimGrid, dimBlock >>>(
                          d_y, d_x, alpha, beta, n_loop);

循环次数小于或等于 24，因为最大递归深度为 24。除了简单的循环操作外，您不必在主机上放置循环操作，因为它已在内核代码中定义。

1. 使用以下命令编译代码：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -rdc=true -o cuda_kernel ./cuda_kernel.cu

您必须使用 GPU 的计算能力版本号来选择gencode选项。

1. 测试编译后的应用程序。这个结果是使用 Tesla P40 测量的，因为 CUDA 9.x 不支持 Volta GPU 的 CUDA 动态并行性（CDP）配置文件：

Elapsed Time...
simple loop: 0.094 ms
inner loop : 0.012 ms
recursion : 0.730 ms

三种执行的比较

从结果中，我们可以确认内部循环是迭代操作中最快的方法。以下截图显示了这个示例应用程序的分析结果：

迭代内核调用显示了每个内核调用的内核启动开销。GPU 需要从设备内存中获取所有所需的数据，并需要调度 GPU 资源等。另一方面，内部循环内核显示了一个打包操作，因为所有所需的资源都是预先定位的，不需要重新调度其执行。由于我们之前讨论的动态并行性限制，递归内核操作显示了最长的执行时间。

一般来说，建议使用开销最小的方法。然而，很难说哪种内核调用设计优于其他，因为算法和问题比我们在这里涵盖的要多。例如，CDP 用于增强某些情况下的并行性，比如用于 GPU 树和搜索。

总结

在本章中，我们涵盖了几种内核执行机制。我们讨论了 CUDA 流是什么，以及如何使用它们同时执行多个内核函数。通过利用主机和 GPU 之间的异步操作，我们学到可以通过数据传输和内核执行来隐藏内核执行时间。此外，我们可以使用回调函数使 CUDA 流调用主机函数。我们可以创建一个有优先级的流，并确认其有优先级的执行。为了测量内核函数的确切执行时间，我们使用了 CUDA 事件，并且我们也学到 CUDA 事件可以用于与主机同步。在最后一节中，我们还讨论了每种内核执行方法的性能。

我们还涵盖了其他内核操作模型：动态并行性和网格级协作组。动态并行性使得内核函数内部可以进行内核调用，因此我们可以使用递归操作。网格级协作组实现了多功能的网格级同步，我们讨论了这个特性在特定领域的用途：图搜索、遗传算法和粒子模拟。

然后，我们扩展了我们对主机的覆盖范围。CUDA 内核可以从多个线程或多个进程中调用。为了执行多个线程，我们使用了带有 CUDA 的 OpenMP，并讨论了它的用处。我们使用 MPI 来模拟多进程操作，并且可以看到 MPS 如何提高整体应用程序性能。

正如我们在本章中看到的，选择正确的内核执行模型是一个重要的话题，线程编程也是如此。这可以优化应用程序的执行时间。现在，我们将扩展我们的讨论到多 GPU 编程来解决大问题。

第五章：CUDA 应用程序分析和调试

CUDA 为开发人员提供了许多编程工具。这些工具包括编译器、分析器、IDE 及其插件、调试器和内存检查器。了解这些工具将有助于您分析您的应用程序，并帮助您完成我们将要涵盖的开发项目。在本章中，我们将介绍这些工具的基本用法，并讨论如何将它们应用到应用程序开发中。

本章将涵盖以下主题：

• 在 GPU 应用程序中进行专注的分析目标范围

• 针对远程机器的可视化分析

• 使用 CUDA 错误调试 CUDA 应用程序

• 使用 CUDA Assert 断言本地 GPU 值

• 使用 Nsight Visual Studio Edition 调试 CUDA 应用程序

• 使用 Nsight Eclipse Edition 调试 CUDA 应用程序

• 使用 CUDA-GDB 调试 CUDA 应用程序

• 使用 CUDA-memcheck 进行运行时验证

技术要求

为了完成本章，建议您使用 Pascal 架构之后的 NVIDIA GPU 卡。换句话说，您的 GPU 的计算能力应该等于或大于 60。如果您不确定您的 GPU 架构，请访问 NVIDIA 的网站developer.nvidia.com/cuda-gpus，并确认您的 GPU 的计算能力。

本章的示例代码已经使用 CUDA Toolkit 的 10.1 版本进行开发和测试。一般来说，如果适用的话，建议您使用最新的 CUDA 版本。

在 GPU 应用程序中进行专注的分析目标范围

NVIDIA 的 Visual Profiler 是一个方便的工具，用于找出 GPU 应用程序中的瓶颈并理解它们的操作。虽然它提供了应用程序操作的流畅信息，但如果您只想专注于特定代码区域，这些信息可能会显得多余。在这种情况下，限制分析范围更加高效。

分析目标可以是特定的代码块、GPU 和时间。指定代码块称为专注分析。当您想要专注于特定内核函数的分析，或者在大型 GPU 应用程序的一部分上进行分析时，这种技术是有用的。在我们介绍专注分析后，将介绍针对 GPU 或时间的分析目标。

限制代码中的分析目标

为了从专注的分析中受益，您可能希望在源代码中包含特色的头文件，如下所示：

#include <cuda_profiler_api.h>

然后，您可以使用cudaProfilerStart()和cudaProfilerStop()来指定您的分析范围：

cudaProfilerStart();
... {target of profile} ...
cudaProfilerStop();

现在，您需要使用特定标志--profile-from-start来分析您的应用程序。

这个选项不会让分析器开始分析，直到请求到达。如果您想使用 NVIDIA Visual Profiler 来分析您的应用程序，请确保在设置视图中勾选“启动时启用分析”复选框。

以下步骤涵盖了如何使用一些简单的示例代码来控制 NVIDIA 分析器。为了使这更容易，我们将重用我们在第三章中用于矩阵乘法操作的示例代码，CUDA 线程编程：

1. 编写一个 CUDA 应用程序，其中包含两个简单的 SGEMM CUDA 内核函数。这两个内核函数是相同的，但名称不同，即sgemm_kernel_A()和sgemm_kernel_B()。

2. 进行两次迭代调用，如下所示：

int n_iter = 5;
for (int i = 0; i < n_iter; i++)
    sgemm_gpu_A(d_A, d_B, d_C, N, M, K, alpha, beta);
for (int i = 0; i < n_iter; i++)
    sgemm_gpu_B(d_A, d_B, d_C, N, M, K, alpha, beta);

1. 现在，让我们编译代码并使用nvprof进行分析：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -o sgemm sgemm.cu
$ nvprof -f -o profile-original.nvvp ./sgemm

当您使用 Visual Profiler 打开生成的profile-original.nvvp文件时，您将得到如下的分析结果：

这个时间轴包括了应用程序启动时的整个分析信息。然而，当我们想要优化我们的内核函数时，我们可以说分析结果包含了不必要的信息。

以下步骤涵盖了如何指定分析专注区域：

1. 在源代码顶部放置 #include <cuda_profiler_api.h> 以启用专注分析 API。然后，我们可以使用 cudaProfilerStart() 和 cudaProfilerStop() 来包含我们感兴趣的区域，如下所示：

cudaProfilerStart();
for (int i = 0; i < n_iter; i++)
    sgemm_gpu_B(d_A, d_B, d_C, N, M, K, alpha, beta);
cudaProfilerStop();

1. 编译您的代码并使用 Visual Profiler 查看更新后的分析结果。我们必须向分析器提供 --profile-from-start off 选项，如下所示：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -o sgemm sgemm.cu
$ nvprof -f -o profile-start-stop.nvvp --profile-from-start off ./sgemm

当您打开新生成的分析结果时，分析器只会报告应用程序的指定部分，如下所示：

分析结果受限。上面的屏幕截图显示了内核执行的情况，从开始 GPU 执行时开始。因此，您可以省去对应用程序初始化和其他无关操作进行分析的步骤。

总之，专注分析有几个好处，如下所示：

• 这有助于您专注于当前正在开发的模块。

• 它可以让您在分析报告中删除无关的操作，例如：

• 与您的代码无关的外部模块行为

• 应用程序初始化延迟

• 在时间轴视图中查找目标函数时，这有助于节省时间。

通过时间或 GPU 限制分析目标

NVIDIA 分析器还有其他可以限制分析目标的选项。您也可以使用以下选项进行专注分析：

• --timeout <second> 选项限制应用程序的执行时间。当您需要分析执行时间较长的迭代操作的应用程序时，此选项非常有用。

• --devices <gpu ids> 选项指定要进行分析的 GPU。该选项帮助您在多 GPU 应用程序中缩小 GPU 内核操作的范围。

此外，如果您只想专注于少数内核函数，您不必收集所有指标。您可以使用 --kernels、--event 和 --metrics 选项向分析器指定您的兴趣。您可以将这些选项与其他分析选项一起使用，如下所示：

$ nvprof -f -o profile_kernels_metric.nvvp --kernels sgemm_kernel_B --metrics all ./sgemm

将收集的指标导入时间轴分析结果后，您会发现目标内核只有指标信息。

在 CPU 抽样中有许多其他多功能的分析特性，例如标记分析范围、OpenMP 和 OpenACC 分析等。如果您想了解 NVIDIA 分析器的功能，请查看 NVIDIA 的 Jeff Larkin 提供的以下分析器介绍讲座：www.olcf.ornl.gov/wp-content/uploads/2018/12/summit_workshop_Profilers.pdf。

NVIDIA 的官方分析器用户指南提供了有关 NVIDIA 分析器功能的详细信息 (docs.nvidia.com/cuda/profiler-users-guide/index.html).

使用 NVTX 进行分析

通过专注分析，我们可以使用 cudaProfilerStart() 和 cudaProfilerStop() 对有限的特定区域进行分析。但是，如果我们想要分析复杂应用程序中的功能性能，这是有限的。对于这种情况，CUDA 分析器通过 NVIDIA 工具扩展 (NVTX) 提供时间轴注释。

使用 NVTX，我们可以对 CUDA 代码进行注释。我们可以使用 NVTX API 如下：

nvtxRangePushA("Annotation");
.. { Range of GPU operations } ..
cudaDeviceSynchronization();     // in case if the target code block is pure kernel calls
nvtxRangePop();

如您所见，我们可以将一段代码定义为一组代码，并手动注释该范围。然后，CUDA 分析器提供注释的时间轴跟踪，以便我们可以测量代码块的执行时间。这种方法的一个缺点是 NVTX API 是主机函数，因此如果目标代码块是纯 GPU 内核调用，则需要同步主机和 GPU。

要了解更多信息，请将此 NVTX 代码应用于前面的专注分析示例。首先，我们应该包含一个 NVTX 头文件，如下所示：

#include "nvToolsExt.h"

然后，我们将在几个地方插入 nvtxRangePushA() 和 nvtxRangePop()，如下所示：

    cudaProfileStart();
    // copy initial value for gpu memory
    nvtxRangePushA("Data Transfer");
    cudaMemcpy(d_A, A, N * K * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, A, K * M * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_C, A, N * M * sizeof(float), cudaMemcpyHostToDevice);
    nvtxRangePop();

    nvtxRangePushA("Kernel Execution");
    // do operation
    nvtxRangePushA("Kernel A");
    for (int i = 0; i < n_iter; i++)
        sgemm_gpu_A(d_A, d_B, d_C, N, M, K, alpha, beta);
    cudaDeviceSynchronize();
    nvtxRangePop();    // Kernel A

    nvtxRangePushA("Kernel B");
    for (int i = 0; i < n_iter; i++)
        sgemm_gpu_B(d_A, d_B, d_C, N, M, K, alpha, beta);
    cudaDeviceSynchronize();

    nvtxRangePop();    // Kernel B
    nvtxRangePop();    // Kernel Execution
    cudaProfileStop();

在上面的代码中，我们已经扩大了关注的配置文件区域，以监视 NVTX 操作。我们还有Data Transfer、Kernel A、Kernel B和Kernel Execution作为 NVTX 范围。NVTX 支持多级注释，因此Kernel A和Kernel B范围将包含在Kernel Execution时间轴中。

要编译代码，我们应该为nvcc编译器提供-lnvToolsExt选项来提供 NVTX API 的定义。我们可以使用以下命令编译代码：

$ nvcc -m64 -gencode arch=compute_70,code=sm_70 -lnvToolsExt -o sgemm sgemm.cu

然后，NVIDIA 分析器可以在没有额外选项的情况下收集 NVTX 注释。我们可以使用以下命令对应用程序进行分析：

$ nvprof -f --profile-from-start off -o sgemm.nvvp ./sgemm.nvvp

以下屏幕截图显示了时间轴分析结果。在这个截图中，我们可以看到用绿色标记的标记和范围。这些绿色条有注释：

前面的屏幕截图为我们提供了以下信息：

• 我们可以通过 NVTX 注释来确定内存复制操作的位置。

• 我们可以通过包装区域来划分功能位置，例如kernel A和kernel B。

• NVTX 注释可以堆叠多个级别的注释。正如我们所看到的，kernel A和kernel B包含在kernel execution注释中。

以下文件不仅介绍了 NVTX，还解释了如何使用 NVTX 来使用不同的颜色：devblogs.nvidia.com/cuda-pro-tip-generate-custom-application-profile-timelines-nvtx。NVTX 的一个应用是使用 NVTX 注释对深度学习网络进行分析。这提供了对网络操作瓶颈的洞察。我们将在本书的第十章《使用 CUDA 进行深度学习加速》中讨论这一点。

针对远程机器进行可视化分析

NVIDIA Visual Profiler 还可以分析远程应用程序。这个功能在远程应用程序开发时特别方便，尤其是在服务器端开发应用程序时。

有几种使用可视化分析器的方法，如下所示：

• 在主机上进行 CUDA 应用程序的分析

• 通过在目标端使用nvprof CLI 收集配置文件数据，将文件复制到主机并使用 Visual Profiler 打开

• 在目标平台上使用主机机器进行应用程序的分析

在主机机器上直接进行可视化分析非常方便，可以节省开发时间。此外，远程分析提供了与在主机机器上分析 GPU 应用程序相同的用户体验。唯一的例外是我们需要建立远程连接。主机管理的可视化分析提供的另一个好处是分析器会自动按需收集度量信息。

NVIDIA 分析器与主机机器中的 NVIDIA 分析器进行通信并收集分析数据。因此，您需要确认您的主机机器（台式机或笔记本电脑）应连接到远程机器。以下图显示了此连接的概述：

让我们尝试远程分析 GPU 应用程序。以下步骤介绍了如何在 NVIDIA Visual Profiler 中分析远程 GPU 应用程序：

1. 首先，转到文件 | 新建会话。当您单击新建会话菜单时，您将看到以下对话框窗口：

1. 然后，我们需要添加一个连接，方法是转到“管理连接...”菜单。然后，将出现“新的远程连接”对话框。通过单击“添加”按钮并在适当的部分输入远程机器信息来添加远程机器信息。然后，通过单击“完成”按钮关闭对话框。完成后，您将看到以下输出：

正如我们之前讨论的，主机和远程机器通过 SSH 进行通信，其默认端口号为 22。如果主机机器使用其他端口进行 SSH，您必须在新的远程会话创建对话框中通知它该端口号。

1. 现在，我们需要通过单击 Toolkit/Script*右侧的“管理...”按钮在远程机器上设置 CUDA Toolkit 路径。*一个很好的开始是使用“检测”按钮。它会自动查找nvcc路径并自动设置配置信息。如果自动检测失败，您必须手动输入配置信息。完成配置过程后，单击“完成”按钮，如下所示：

1. 通过单击“浏览”按钮在“文件”文本框的右侧指定 GPU 应用程序的二进制文件。它会要求您的远程机器登录密码。找到应用程序路径并设置应用程序路径。如果需要控制应用程序的行为，还可以输入应用程序的参数。完成应用程序和连接设置后，单击“下一步”按钮设置分析器的选项。

2. 现在，我们将设置分析器选项。NVIDIA Visual Profiler 允许我们使用复选框设置分析器的选项，如下面的屏幕截图所示。单击“完成”，分析器将从应用程序收集分析数据：

您将在主机机器上看到时间线分析输出。

1. 最后，分析分析时间线图的性能。单击要分析的任何内核函数。单击“执行内核分析”按钮；分析工具将收集相关的度量信息。通过这样做，您可以快速获得有关性能限制器的报告，并找到内核函数的瓶颈。

使用 CUDA 错误调试 CUDA 应用程序

具有专用的异常检查和检查错误是使软件具有高质量的基本特征之一。CUDA 函数通过返回每个函数调用的状态来报告错误。不仅如此，CUDA API，而且内核函数和 CUDA 库的 API 调用也遵循这个规则。因此，检测到重复错误是识别 CUDA 执行中错误的开始。例如，假设我们使用cudaMalloc()函数分配了全局内存，如下所示：

cudaMalloc((void**)&ptr, byte_size);

如果全局内存没有足够的空闲空间来分配新的内存空间怎么办？在这种情况下，cudaMalloc()函数返回一个错误来报告内存不足异常。通过使用cudaGetLastError()可以捕获由内核调用触发的标志。它返回记录的错误状态并重置标志的值。但是要小心处理这个标志：它的返回并不保证错误发生在 GPU 的最后执行，并且需要手动重置标志。

CUDA API 的返回值和cudaGetLastError()函数的返回值都是cudaError_t类型。这种cudaError_t类型是预定义的整数类型，应用程序可以识别发生了哪种类型的错误。例如，此类型定义如下：

Enum cudaErorr_t {
    cudaSuccess = 0,
    cudaErrorMemoryAllocation = 2, 
    cudaErrorUnknown = 30,
    cudaErrorNoDevice = 38,
    cudaErrorAssert = 59,
    cudaErrorTooManyPeers = 60,
    cudaErrorNotSupported = 71,
    ....
};

记住或翻译所有这些值是不切实际的。为此，CUDA 示例代码提供了一个辅助函数checkCudaError()，它位于common/inc/cuda_helper.h中。当 CUDA 函数返回错误时，此函数打印出错误消息。其函数定义如下：

#define checkCudaErrors(err) { \
    if (err != cudaSuccess) {  \

        fprintf(stderr, "checkCudaErrors() API error = %04d \"%s\" from file <%s>, line %i.\n", \
                err, cudaGetErrorString(err), __FILE__, __LINE__); \
        exit(-1); \
    } \
}
#endif

由于此函数被定义为宏，我们可以确定发生错误的行。

我们可以使用这个函数的两种方式。一种是在源代码中包含cuda_helper.h文件。另一种是将函数代码复制到代码中的某个位置。

然后，我们将使用checkCudaErrors()包装所有的 CUDA API 类，如下所示：

checkCudaErrors(cudaMalloc((void **)&d_A, N * K * sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_B, K * M * sizeof(float)));
checkCudaErrors(cudaMalloc((void **)&d_C, N * M * sizeof(float)));

对于内核函数调用，我们将使用cudaGetLastError()函数来获取内核调用的错误标志，如下所示：

sgemm_kernel_A<<<dimGrid, dimBlock>>>(A, B, C, N, M, K, alpha, beta);
checkCudaErrors(cudaGetLastError());

然而，这段代码有一个问题：内核操作与主机异步，所以cudaGetLastError()只能捕获主机端的返回值。很可能错误是在应用程序的某个地方触发的。为了解决这种情况，您可以使用任何主机和设备同步函数；例如：

sgemm_kernel_A<<<dimGrid, dimBlock>>>(A, B, C, N, M, K, alpha, beta);
checkCudaErrors(cudaDeviceSynchronize());

现在，让我们通过修改源代码来测试错误检测代码。例如，您可以请求cudaMemcpy复制比分配大小更大的内存空间。在这种情况下，应用程序会返回一个错误消息，如下所示：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -lnvToolsExt -o sgemm ./sgemm.cu
CUDA error at sgemm.cu:93 code=11(cudaErrorInvalidValue) "cudaMemcpy(d_A, A, N * K * sizeof(float), cudaMemcpyHostToDevice)"

或者，您可以为 CUDA 内核传递一个NULL指针，以便内核访问无效的内存空间。在这种情况下，应用程序会在cudaDeviceSynchronize()中报告非法地址错误：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -lnvToolsExt -o sgemm ./sgemm.cu
CUDA error at sgemm.cu:104 code=77(cudaErrorIllegalAddress) "cudaDeviceSynchronize()"

这个错误检查宏非常有用，因为它报告了错误发生的源代码位置。然而，这个报告有一个缺点，就是它检测到的错误位置与实际发生的错误位置不匹配。

错误消息应该报告我们复制比分配的内存更大的内存位置导致非法值错误。因此，开发人员可以在内核调用之后立即识别错误消息。然而，这个错误检查代码只在主机上工作。因此，如果 GPU 操作没有正确同步，这可能会混淆 GPU 操作。例如，如果我们没有设置同步，只是检查错误，那么cudaDeviceSynchronize()函数可能会报告错误的位置。在这种情况下，我们可以设置CUDA_LAUNCH_BLOCKING=1环境变量，使所有内核执行与主机同步：

$ ./sgemm
CUDA error at sgemm.cu:104 code=77(cudaErrorIllegalAddress) "cudaDeviceSynchronize()" 
$ CUDA_LAUNCH_BLOCKING=1 ./sgemm
CUDA error at sgemm.cu:36 code=77(cudaErrorIllegalAddress) "cudaGetLastError()"

sgemm.cu的第 36 行是cudaGetLastError()调用，在sgemm内核调用之后。这就是我们放置一个预期错误的位置。我们可以在运行时确定正确的错误位置。

有两份官方文件可以帮助您了解不同类型的 CUDA 错误：

• docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html

• CUDA Toolkit 根路径中的include/driver_types.h

使用 CUDA assert 断言本地 GPU 值

即使您的 GPU 应用程序没有任何系统错误，您也需要检查计算结果，以确保执行的结果符合设计要求。为此，CUDA 提供了assert函数，它检查参数值是否为零。如果是，这个函数会引发一个错误标志，以便主机可以识别内核函数中存在错误。

断言用于验证操作结果是否符合预期。在 CUDA 编程中，可以从设备代码中调用assert函数，并在给定参数为零时停止内核的执行：

void assert(int expression);

这是assert函数的声明，与 C/C++的声明相同。当断言被触发时，应用程序会停止并报告其错误消息。如果应用程序由调试器启动，它会作为断点工作，以便开发人员可以调试给定的信息。例如，输出消息看起来像这样：

$ nvcc -run -m64 -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -lnvToolsExt -o sgemm ./sgemm.cu
sgemm.cu:29: void sgemm_kernel_A(const float *, const float *, float *, int, int, int, float, float): block: [16,64,0], thread: [0,0,0] Assertion `sum == 0.f` failed.     

由于输出消息指向了确切的 CUDA 块和线程索引，开发人员可以轻松分析指定的 CUDA 线程的执行。

现在，让我们应用断言并看看它如何检测到预期的错误。我们将修改在GPU 应用程序中的性能优化目标范围部分中使用的 SGEMM 操作代码。

首先，在内核函数的中间放置断言代码。我们将看到表达式的效果，它应该是 false。断言代码可以编写如下：

__global__ void sgemm_kernel_A(const float *A, const float *B, float *C, int N, int M, int K, float alpha, float beta)
  {
      int col = blockIdx.x * blockDim.x + threadIdx.x;
      int row = blockIdx.y * blockDim.y + threadIdx.y;
      float sum = 0.f;
      for (int i = 0; i < K; ++i) 
          sum += A[row * K + i] * B[i * K + col];

      if (row == 0 && col == 0)
 assert(sum == 0.f);

      C[row * M + col] = alpha * sum + beta * C[row * M + col];
  }

您可以尝试其他索引值或尝试其他可能的错误。编译代码并运行它以查看输出。以下代码显示了此修改的输出错误：

sgemm.cu:29: void sgemm_kernel_A(const float *, const float *, float *, int, int, int, float, float): block: [0,0,0], thread: [0,0,0] Assertion `sum == 0.f` failed.

错误消息报告了断言触发的代码位置、内核函数的名称和 GPU 的线程索引。有了这些信息，我们可以很容易地找出应该从哪里开始分析。

实际上，assert函数的使用与普通 C/C++编程中的assert函数相同。一个区别是assert函数在设备代码中起作用。因此，它不仅报告事件位置和表达式，还显示块和线程索引。

然而，使用断言会对应用程序性能产生影响。因此，我们应该只在调试目的时使用断言。建议在生产环境中运行时禁用它。您可以通过在包含assert.h之前添加NDEBUG预处理宏来在编译时禁用断言。

使用 Nsight Visual Studio Edition 调试 CUDA 应用程序

对于 Windows 应用程序开发人员，CUDA Toolkit 提供了 Nsight Visual Studio Edition，它可以在 Visual Studio 中实现 GPU 计算。这个工具作为 Visual Studio 的扩展工作，但您可以构建、调试、分析和跟踪 GPU 应用程序以及主机。如果您的工作平台不是 Windows，则本节中的内容将不适用，您可以跳过它。

CUDA 调试器允许我们监视每个 CUDA 线程的 GPU 内核上的本地值。与普通主机调试一样，您可以在内核代码中设置断点并触发它们。您还可以设置条件，例如其他普通断点。有了这个功能，您可以为特定的 CUDA 线程索引触发断点并查看它们的本地变量。

这个工具可以与 CUDA Toolkit 一起安装。您可以从网站上获取最新版本。这不是强制性的，但是当您的开发环境使用旧的 CUDA Toolkit 和最新的 GPU 及其驱动程序时，建议使用它。访问 NVIDIA Nsight 网页（developer.nvidia.com/nsight-visual-studio-edition）下载并安装 Nsight。您需要 NVIDIA 开发人员会员资格才能获取该软件。您还需要安装推荐的显示驱动程序版本。

您可以通过转到 Visual Studio 菜单栏中的菜单 | Nsight 来找到 CUDA 工具。此菜单中有几个工具，其中一些如下：

• 图形调试：用于图形（Direct3D、OpenGL 和 Vulkan）应用程序的调试器

• CUDA 调试（Next-Gen）：用于同时调试 CPU 和 GPU 代码的调试器（Turing、Volta 和 Pascal 与最新驱动程序）

• CUDA 调试（传统）：仅用于 GPU 内核的调试器（具有旧驱动程序的 Pascal、Maxwell 和 Kepler）

• 性能分析：用于分析当前 GPU 应用程序的性能

• CUDA 内存检查器：用于在运行时检查 GPU 内存违规（如前一节中介绍的）

在本节中，我们将重点放在 CUDA 调试（Next-Gen）上。这是因为 Next-Gen 调试器可以支持包括 Turing 和 Volta 在内的最新架构。CUDA 内存检查器将在本章末尾介绍。

现在，让我们配置一个示例项目，并看看我们如何使用 Nsight Visual Studio Edition 调试应用程序。您可以使用默认的示例代码，或者用我们之前介绍的 CUDA 代码替换代码。您还可以使用05_debug/05_debug_with_vs文件中提供的示例代码。这是一些简单的 SAXPY 代码。

将项目属性设置为生成适当的设备目标代码。在项目的属性页面中，您可以指定目标代码版本。在 CUDA C/C++ | 代码生成文本框中列出您想要在其中使用的架构版本：

上述截图显示了 CUDA 设备代码生成属性页面。您可以设置几个nvcc选项，例如目标 GPU 的计算能力、每个线程的寄存器限制以及在编译时冗长的 CUDA 内核信息。

在第 34 行和第 75 行设置断点，其中第 34 行是内核函数的中间位置，第 75 行是从主机复制数据到设备的位置。然后，使用以下方法之一编译并开始调试：

• 在 Visual Studio 菜单栏中导航到 Nsight，然后单击“开始 CUDA 调试（Next-Gen）”。

• 在“解决方案资源管理器”中右键单击项目，选择“调试|开始 CUDA 调试（Next-Gen）”。

• 转到 Nsight CUDA 调试工具栏，单击“开始 CUDA 调试（Next-Gen）”。

Windows 防火墙可能会询问您是否信任并允许 Nsight 的网络连接。这是正常的，因为 Nsight 使用内部网络来监视 GPU 设备。单击“接受”并继续调试。当前的 Nsight Visual Studio Edition 提供了两种调试选项。这取决于目标 GPU 架构版本。如果您的 GPU 是 Volta 或 Turing，建议使用“Next-Gen”调试。如果您的 GPU 是 Pascal，则适当的调试器取决于驱动程序版本。为了澄清，请访问 NVIDIA 支持的 GPU 列表：developer.nvidia.com/nsight-visual-studio-edition-supported-gpus-full-list。

应用程序将在应用程序启动的地方停止。继续跟踪。应用程序将在主机的第 75 行和设备的第 34 行停止。从中我们可以了解到，Nsight 可以同时跟踪主机和设备上的 GPU 应用程序。

当黄色箭头停在内核函数中时，您可以查看局部变量。全局索引中的线程索引为0。由于 CUDA 并行发出多个 CUDA warp 和 CUDA 线程，因此您可以通过更改blockIdx和threadIdx来查看其他线程的局部变量。基本的 CUDA 线程调试控制单元是 warp。换句话说，您可以控制调试器以遍历活动 warp。Nsight 调试器在 Nsight 菜单栏中的“上一个活动 warp/下一个活动 warp”菜单中提供了此功能。

以下屏幕截图显示了我们在调试时出现的 Nsight 调试控件：

如果更改 warp，您会发现在“Autos”面板中监视的局部变量会随着 warp 的变化而更新索引。例如，以下屏幕截图显示了“Autos”窗口，该窗口报告了活动 warp 中所选线程的局部变量，即正在由主导线程监视的局部变量的值：

选择的线程更改后，Autos 值会更新。以下屏幕截图显示了通过移动到下一个活动 warp 所做的更改：

Next-Gen CUDA 调试器提供了三种类型的窗口——warp info、lanes 和 GPU registers。黄色箭头表示当前的 GPU 执行，并以三个方面显示其信息：

• Warp Info 窗口提供了另一种选择活动 warp 的方法。您可以在菜单栏中从 Nsight | Window | Warp Info 打开该窗口。窗口如下所示：

每行表示 CUDA 网格中的活动 warp。第四列“Shader Info”显示了每个 warp 的块和主导线程索引。第五列“threads”显示了 warp 中 CUDA 线程的状态。单元格的颜色表示每个线程的状态。由于我们在断点处观察它们，它们都是红色的，但在调试过程中您会看到其他颜色。以下屏幕截图解释了每种颜色在线程状态方面的含义：

双击任何 warp，查看 autos 窗口中的局部变量是如何更新的。

• Lanes 窗口允许您在所选活动 warp 内选择特定的 CUDA 线程。一个 lane 指的是 warp 中的一个线程。您可以从 Nsight | Window | Lanes 中打开该窗口。通过双击一个 lane，您可以发现 autos 窗口中的局部变量根据更新的索引而更新：

活动 warp 中的 lanes 窗口信息。

寄存器窗口显示了 GPU 寄存器的当前状态。如果它们的值被更新，它们将变为红色。

如果您想了解如何使用 Nsight Visual Studio Edition，请阅读 NVIDIA 官方用户指南。它介绍了如何配置调试环境，如何使用它，以及各种情况下的详细提示。 (docs.nvidia.com/nsight-visual-studio-edition/Nsight_Visual_Studio_Edition_User_Guide.htm)。

使用 Nsight Eclipse Edition 调试 CUDA 应用程序

对于 Linux 和 OSX 平台开发，CUDA Toolkit 提供了 Nsight Eclipse Edition。这个工具基于 Eclipse，因此开发人员可以很容易地在 CUDA C 开发中使用这个工具。

Nsight Eclipse Edition 是基于 Eclipse 用于 CUDA 应用程序开发的。您可以使用它来编辑、构建、调试和分析 CUDA 应用程序。它使得在 Linux 和 OSX 中进行 CUDA C/C++开发变得简单。这个工具作为 CUDA Toolkit 的一部分安装，因此您不必单独安装这个工具。但是，如果您使用 Linux，需要配置 Java 7 才能使用它。

Nsight Eclipse Edition 是基于 Eclipse 4.4.0 版本（2014 年发布的 Luna 版本）构建的，并且基于 Java 7 构建。

Nsight 可以通过终端中的nsight命令或者 X 窗口应用程序列表中执行。

现在，让我们从终端或 X 窗口桌面打开 Nsight，以便我们可以编译和分析给定的示例。要么创建一个新的 CUDA 项目，要么打开05_debug/06_debug_with_eclipse中提供的示例项目。如果要创建项目，请选择 CUDA C/C++项目。空项目只会给您一个空项目，而 CUDA Runtime 项目会给您一个带有一些示例代码的项目。如果要使用示例项目，请使用文件 | 导入 | 导入现有项目到工作区。

现在，让我们在sgemm内核函数中设置一个断点。就像在 Eclipse 中的普通 C/C++项目一样，您可以在nsight中构建和调试 CUDA 应用程序。在内核函数的起始点（第 23 行）设置一个断点，如下所示：

对于内核函数调试来说，一个很好的起点是在线程索引计算之后。设置一个断点来暂停 GPU 的执行。现在，通过单击菜单面板中的绿色 bug 来编译和开始调试。在调试窗口切换调试透视之时，点击继续，直到达到我们设置的断点。

Nsight 允许您监视活动 warp 中的局部变量和寄存器。首先，它会在 CUDA 网格中的领先 CUDA 线程（CUDA 线程0）处停止应用程序。然后，您可以从调试窗口切换到其他 CUDA 活动 warp，并使用 CUDA 窗口检查每个 CUDA 线程，就像这样：

以下截图显示了所选 CUDA 线程的局部变量信息。Nsight 会在这些值更新时更新它们：

上述截图显示了 Eclipse 的调试透视窗口中的 Debug 窗口和 CUDA 窗口。调试窗口提供了在所选 GPU 上的活动 warp 中进行 CUDA warp 选择的功能，并且可以在所选活动 warp 内进行 lane 选择。

NVIDIA 还有一个 Nsight Eclipse Edition 用户指南。您可以通过访问docs.nvidia.com/cuda/nsight-eclipse-edition-getting-started-guide/index.html来了解更多关于这个工具的信息。

使用 CUDA-GDB 调试 CUDA 应用程序

CUDA 工具包提供了 CUDA-GDB，它支持 CUDA C/C++调试，用于诸如 C/C++ GDB 之类的程序。这对于直接调试没有 X 窗口环境或远程调试的 CUDA C/C++应用程序非常有用。

要调试 GPU 应用程序，Makefile应该包括主机的-g调试标志和 GPU 的-G调试标志。基本上，CUDA 的 GDB 用法与主机调试相同，只是在 CUDA 操作之外还有一些额外的调试功能。例如，我们可以设置特定的 CUDA 线程和 CUDA 感知断点。

CUDA-GDB 的断点

让我们看看cuda-gdb如何帮助我们检测代码中的错误。我们将在代码中设置断点，并查看主机和 GPU 上的局部值。为此，将工作目录切换到05_debug/07_debug_with_gdb目录。我们将通过将其与适当的行匹配来检查cuda-gdb的操作。

首先，让我们使用以下命令编译源代码：

$ nvcc -run -m64 -g -G -Xcompiler -rdynamic -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o simple_sgemm ./simple_sgemm.cu

然后，我们应该执行cuda-gdb，这样我们就可以在终端上调试应用程序，如下所示：

$ cuda-gdb simple_sgemm

我们可以在代码的特定行上设置断点，如下所示：

(cuda-gdb) break simple_gemm.cu:21

或者，我们可以按照内核函数的名称设置断点，如下所示。这将在函数的入口点触发断点：

(cuda-gdb) break sgemm_kernel

如果cuda-gdb警告指出断点希望在未来的共享库加载时挂起，则回答y。您也可以在主机代码上设置断点。

使用断点的一个问题是，断点将根据 CUDA 线程的数量触发。因此，我们应该提供条件信息，以便针对特定的 CUDA 线程设置断点。条件断点如下：

(cuda-gdb) break sgemm_kernel if blockIdx.y == 2

当然，我们可以修改预定义断点的条件如下：

(cuda-gdb) cond 3 // break 3 is defined previously

让我们使用run命令执行示例应用程序。如果应用程序遇到任何断点，CUDA-GDB 将提供有关它的信息。以下代码显示了应用程序在第21行遇到断点时cuda-gdb的报告：

(cuda-gdb) run
[Switching focus to CUDA kernel 0, grid 1, block (0,0,0), thread (5,0,0), device 0, sm 0, warp 0, lane 5]
Thread 1 "simple_sgemm" hit Breakpoint 1, sgemm_kernel<<<(128,128,1),(16,16,1)>>> (A=0x7fffb6000000, B=0x7fffb7000000, C=0x7fffb4000000, N=2048, M=2048, K=2048, alpha=2, beta=1) at simple_sgemm.cu:21
21 int col = blockIdx.x * blockDim.x + threadIdx.x;

现在，是时候使用 GDB 命令来跟踪代码或监视活动变量了。我们可以使用 next（或n）、step（或s）、continue（或c）和 finish（或fin）来跟踪内核函数。然而，当我们到达内核代码的末尾并需要在主机和设备之间切换目标硬件时，我们应该使用continue命令。

使用 CUDA-GDB 检查变量

除了默认的 GDB 命令之外，CUDA-GDB 提供了可以与 CUDA 内核一起使用的调试功能。以下是您可以使用 CUDA-GDB 做的事情。

列出内核函数

与普通函数一样，CUDA-GDB 可以在内核函数上设置断点。一旦应用程序被断点停止，您可以列出它们如下：

(cuda-gdb) info cuda kernels
Kernel Parent Dev Grid Status   SMs Mask     GridDim  BlockDim Invocation
*      0      -   0    1 Active 0xffffffff (128,128,1) (16,16,1) sgemm_kernel(A=0x7ffff5a79010, B=0x7ffff4a78010, C=0x7ffff3a77010, N=2048, M=2048, K=2048, alpha=2, beta=1)

正如您所看到的，前面的输出显示了内核的配置信息和输入参数变量。

变量调查

CUDA-GDB 帮助我们通过选择特定的线程块索引和线程索引来跟踪特定的 CUDA 线程。有了这个功能，您可以将当前焦点移动到指定的线程。在这个例子中，块大小为 16，col变量被定义为x维度上的 CUDA 线程索引。以下代码显示了 CUDA-GDB 如何通过更改线程索引来报告所选的局部变量的值：

(cuda-gdb) print col
$1 = <optimized out>
(cuda-gdb) cuda kernel 0 block 1,2,0 thread 3,4,0
21 int col = blockIdx.x * blockDim.x + threadIdx.x;
(cuda-gdb) s
22 int row = blockIdx.y * blockDim.y + threadIdx.y;
(cuda-gdb) p col
$2 = 19

检查当前焦点线程的信息：

(cuda-gdb) cuda device kernel block thread
kernel 3, block (1,2,0), thread (3,4,0), device 0

有了手头的信息，我们可以追踪 CUDA 线程。

如果您想了解有关 CUDA-GDB 的更多信息，请查看 NVIDIA 的用户指南文档：docs.nvidia.com/cuda/cuda-gdb/index.html。

使用 CUDA-memcheck 进行运行时验证

CUDA 编程的一个困难点是处理内存空间。由于 CUDA 线程并行操作，边界条件或意外的索引操作可能会违反有效的内存空间。CUDA memcheck 是一个运行时测试工具，如果任何 GPU 操作超出了无效的内存空间，它将验证内存访问。该工具检测以下内存错误：

名称	位置	描述	精确
内存访问错误	设备	无效的内存访问（超出边界，未对齐）	O
硬件异常	设备	硬件错误	X
Malloc/free 错误	设备	在 CUDA 内核中不正确使用malloc()/free()	O
CUDA API 错误	主机	CUDA API 的错误返回	O
cudaMalloc 内存泄漏	主机	使用cudaMalloc()分配的设备内存未被应用程序释放	O
设备堆内存泄漏	设备	在设备代码中使用malloc()分配的设备内存未被应用程序释放	X

精确（O）表示 memcheck 可以指定崩溃的行和文件。另一方面，不精确（X）表示该工具可以识别错误，但由于并发状态，无法指定错误点。cuda-memcheck不需要重新编译进行测试。但是，如果我们使用一些额外的nvcc选项进行编译，我们可以跟踪错误点。nvcc选项包括生成行号信息的-lineinfo和用于保留函数符号的-Xcompiler -rdynamic。

基本上，cuda-memcheck是一个独立的工具，可以在运行时验证 GPU 应用程序。以下命令显示了它在独立模式下的格式：

$ cuda-memcheck [options] <application>

这个工具也可以与 CUDA-GDB 一起使用，帮助开发人员识别错误并进行调试。在 CUDA-GDB 命令行中，使用set cuda memcheck on命令启用内存检查。这样，CUDA-GDB 可以识别与内存相关的异常。

检测内存越界

现在，让我们看看cuda-memcheck如何检测内存异常并与 CUDA-GDB 一起工作。为了简化这个过程，我们将编写一些错误的代码，并查看cuda-memcheck如何报告结果。让我们从一些干净的代码开始。您可以使用05_debug/08_cuda_memcheck中提供的示例代码进行测试。让我们使用cuda-memcheck测试代码并验证它：

$ nvcc -m64 -g -G -Xcompiler -rdynamic -gencode arch=compute_70,code=sm_70 -I/usr/local/cuda/samples/common/inc -o simple_sgemm ./simple_sgemm.cu
$ cuda-memcheck simple_sgemm
========= CUDA-MEMCHECK
Application finished successfully.========= ERROR SUMMARY: 0 errors

现在，让我们将一些错误的代码放入内核函数中，如下所示。如果您愿意，您也可以放入其他错误：

For instance, you may add one to the row value.
__global__ void sgemm_kernel(const float *A, const float *B, float *C, int N, int M, int K, float alpha, float beta)
{
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    row += 1;

    float sum = 0.f;
    for (int i = 0; i < K; ++i)
        sum += A[row * K + i] * B[i * K + col];
    C[row * M + col] = alpha * sum + beta * C[row * M + col];
}

让我们编译并启动代码。内核将返回一个 CUDA 错误，checkCudaErrors()将报告一个错误消息，如下所示：

CUDA error at simple_sgemm_oob.cu:78 code=77(cudaErrorIllegalAddress) "cudaDeviceSynchronize()"

然而，如果我们希望确定内核代码中的哪一行是问题的根本原因，这些信息是不够的。使用cuda-memcheck，我们可以确定哪个 CUDA 线程和内存空间触发了错误，并给出堆栈地址：

$ cuda-memcheck simple_sgemm_oob

输出如下：

前面的屏幕截图显示了cuda-memcheck独立执行的一部分，显示了内核中检测到的所有错误。在这种情况下，cuda-memcheck报告检测到在第 27 行发生的内存违规错误。默认情况下，cuda-memcheck在检测到错误时会停止应用程序的执行。

在这种情况下，我们可以通过检查相关变量来轻松找到根本原因，使用cuda-gdb。为此，我们需要使用cuda-gdb启动应用程序，并启用cuda-memcheck，如下所示：

$ cuda-gdb simple_sgemm_oob
(cuda-gdb) set cuda memcheck on
(cuda-gdb) run

这个过程使cuda-gdb从cuda-memcheck报告非法内存访问检测：

前面的屏幕截图显示了cuda-gdb与cuda-memcheck的报告。开发人员可以轻松地确定simple_sgemm_oob.cu中的第 27 行触发了报告的错误。根据给定的信息，我们可以开始调查哪一块内存访问了无效的空间，如下所示：

(cuda-gdb) print A[row * K + i]
Error: Failed to read generic memory at address 0x7fffc7600000 on device 0 sm 41 warp 20 lane 16, error=CUDBG_ERROR_INVALID_MEMORY_SEGMENT(0x7).
(cuda-gdb) print row * K + i
$1 = 4194304

在不费吹灰之力的情况下，我们可以确定访问A[row * K + i]会触发错误，并且请求的值超出了全局内存（A）的分配空间。通过这种方式，您可以轻松地缩小根本原因。

检测其他内存错误

CUDA memcheck 工具提供了额外的软件验证功能，其中一些如下：

名称	描述	选项
内存泄漏	用于识别内存泄漏	--leak-check full
竞争检查	用于分析多个线程之间对共享内存的冲突访问的竞争危险	--tool racecheck
初始化检查	在没有初始化的情况下识别设备全局内存访问	--tool initcheck
同步检查	验证同步原语的正确使用，如__syncthreads()，__syncwarp()和协作组 API	--tool synccheck

这些工具假设内存访问是正确的或经过验证的，并且不检查内存错误。因此，您需要确认您的应用程序中不存在内存错误。其他有用的 memcheck 选项包括--save，我们可以用它来将输出保存到磁盘，以及--print-level，我们可以用它来控制输出的详细级别。

NVIDIA 为cuda-memcheck提供了用户指南。该文档将帮助您使用 GPU 验证您的应用程序并检测意外错误 (docs.nvidia.com/cuda/cuda-memcheck/index.html)。

使用 Nsight Systems 对 GPU 应用程序进行分析

在本节中，我们将介绍新引入的 CUDA 分析器工具，即 Nsys 和 Nvprof。这些分析器支持 Volta 架构及更高版本的 GPU。它是图灵架构 GPU 中的主要分析器。我们将先介绍 Nsys，然后在下一节介绍 Nvprof。

Nsight Systems (developer.nvidia.com/nsight-systems)是一个系统范围的性能分析工具，可以在时间轴上可视化操作并轻松找到优化点。在时间轴分析方面，Nsight Systems 提供了系统利用率信息，以便我们可以分析瓶颈点。我们可以从 NVIDIA 网站获取 Nsight Systems，但 CUDA 10 默认包含了 Nsight Systems 在工具包中。我们只需要确保它安装正确即可。

对于 CLI，我们应该设置PATH以便于我们的操作，因为它的路径与普通的 CUDA 二进制文件分开。我们可以使用以下命令将其包含在PATH环境变量中：

export PATH=$PATH:/usr/local/cuda/bin:/usr/local/cuda-10.1/NsightSystems-2019.3/Target-x86_64/x86_64

Nsys 提供了两个接口：一个用于 GUI，一个用于 CLI。在主机上，我们可以通过 GUI 运行应用程序来收集应用程序的采样信息。在远程机器上，我们可以通过 CLI 收集分析数据，使用以下命令：

$ nsys profile -t osrt,cuda,nvtx,cublas,cudnn -o baseline -w true <command>

这个选项可以解释如下：

	选项	开关
跟踪	-t/--trace	cuda: 用于跟踪 CUDA 操作，nvtx: 用于跟踪nvtx标签，cublas, cudnn, opengl,openacc: 用于跟踪 API 操作，osrt: 用于跟踪 OS 运行时库，none: 不进行 API 跟踪
输出文件	-o/--output	输出文件名
显示输出	-w/--show-output	true/false: 在终端上打印出分析器的行为

例如，我们可以从02_nvtx SGEMM 应用程序中获得一个名为sgemm.qdrep的分析文件。让我们比较 Nsight Systems 和 NVIDIA Visual Profiler 之间的分析输出。我们可以使用以下命令收集 Nsys 的分析数据：

$ nsys profile -t osrt,cuda,nvtx -o sgemm -w true ./sgemm

这是来自 Nsys 的分析时间轴视图：

以下屏幕截图显示了来自 NVIDIA Visual Profiler 的分析时间轴视图：

Visual Profiler 显示操作事件块，而 Nsight Systems 同时显示系统利用率。因此，我们可以轻松地看到哪些资源（CPU 核心、GPU 或 PCIe 总线）对性能产生影响。此外，Nsight Systems 提供了更具交互性的性能分析体验。当双击任何函数操作时，Nsight Systems Viewer 会展开时间轴以适应窗口，并帮助我们检查操作。此外，Nsight Systems 使我们能够轻松地发现在某个 NVTX 区域下发生的内核执行次数。在 Visual Profiler 时间轴视图中，内核执行看起来像是单个执行，但 Nsight Systems 显示了分离的执行。

现在我们已经确定了应该优化的函数，我们可以继续使用 Nsight Compute，这是另一个新的性能分析器，用于检查内核函数的 GPU 操作。

使用 Nsight Compute 进行内核性能分析

Nsight Compute 是一个用于计算的内核级性能分析器。它收集 GPU 指标信息，并帮助我们专注于 CUDA 内核的优化。换句话说，这个工具涵盖了 Visual Profiler 的性能分析功能。

Nsight Compute 提供两种接口：GUI 和 CLI。GUI 支持主机和远程应用程序性能分析，而 CLI 适用于目标机器。然而，我们可以使用 GUI 获取分析数据并查看结果。

使用 CLI 进行性能分析

为了方便使用 Nsight Compute CLI，我们需要在/usr/local/cuda-10.1/NsightCompute-2019.3/nv-nsight-cu-cli中设置PATH环境变量。然后，我们可以使用以下命令收集性能分析数据：

$ nv-nsight-cu-cli -o <output filename> <application command>

这个命令收集 GPU 执行指标信息，并将数据保存到指定的文件中。如果我们没有提供输出文件名，Nsight Compute 将把收集到的指标报告输出到控制台，从而在控制台上提供快速的指标性能报告。

由于我们可以指定性能分析目标，我们可以限制 Nsight Compute 收集以下信息：

• --kernel-regex：指定要进行性能分析的内核

• --设备：专注于对特定 GPU 进行性能分析

当我们需要在控制台上查看报告时，这个功能非常有用。

使用 GUI 进行性能分析

通过在 Nsight Compute 中打开一个新项目，我们可以启动性能分析操作。以下截图显示了性能分析配置。对于主机应用程序开发，请连接到本地主机。或者，您可以指定要进行性能分析的目标 GPU 服务器：

当然，我们也可以打开使用 CLI 工具在目标机器上生成的nsight-cuprof-report文件。例如，我们可以使用以下命令创建 sgemm 性能分析文件：

$ nv-nsight-cu-cli -o reduction reduction

对于 OSX 用户，Nsight Systems 将需要目标glib库进行远程性能分析。在这种情况下，我们应该从 Nsight Compute 安装映像中复制该库。它将所需的库提供为一个名为 target 的目录，并将该目录复制到Applications/NVIDIA Nsight Compute.app/target目录。

为了方便起见，我们将使用来自第三章 CUDA 线程编程的减少示例代码。它有两个不同寻址的并行减少实现。您可以在03_cuda_thread_programming/05_warp_divergence目录中找到代码。完成连接和应用程序可执行文本栏的设置后，单击启动按钮。然后，按下Ctrl + I，Ctrl + K键以运行到下一个内核函数，然后性能分析器将停在reduction_kernel_1处。按下Ctrl + I，Ctrl + P键以对此内核进行性能分析。然后您将得到以下输出。这张图片展示了 Nsight Compute 基于 GUI 的第一个内核性能分析：

显示基于 GUI 的配置文件（用于第一个内核配置文件）

它提供了交互式配置文件和调试。使用步骤控制调试按钮，我们可以调试 CUDA API 和内核函数。我们还可以使用左侧 API 流面板上的控制按钮移动到下一个内核函数或下一个配置文件范围。在右侧面板上，您可以获取内核的详细配置文件信息。

我们还可以通过启用自动配置文件来自动获取配置文件结果，具体操作如下：转到菜单栏，选择 Profile | Auto Profile。然后，继续进行应用程序。Nsight Systems 将配置所有的内核函数。或者，您可以通过单击窗口顶部的 Profile Kernel 按钮来手动配置内核函数。当我们使用 CLI 收集的配置文件结果时，我们将只看到所有内核函数的配置文件数据。

性能分析报告

在交互式配置文件窗口的右侧面板上，我们可以看到 Nsight Compute 提供了性能分析报告。从报告中，我们可以确定性能限制因素并调查未充分利用的资源。此外，Nsight Compute 还根据资源利用统计数据提供优化建议。我们也可以直接从直接配置文件中识别它们。

此外，Nsight Compute 通过分析 GPU 组件的利用率提供优化建议。它找到瓶颈并建议进行推荐的调查以优化内核。

此报告页面提供了每个组件的利用率，如计算、内存、调度器、指令、warp 等。此外，您可以通过扩展每个组件的左上箭头来获取更多详细信息。以下图片显示了内存工作负载分析的示例报告：

在 Nsight Compute 中，我们可以轻松获取这样的详细信息。在以前的分析器 NVIDIA Profiler 中，我们应该执行每个分析以获取这样的信息。

基线比较

在优化过程中，我们应该将新结果与基线操作进行比较。为了使这项任务对我们来说更容易，Nsight Compute 提供了基线比较功能。单击性能报告面板顶部的 Add baseline 按钮，并将其更改为其他内核函数。然后，我们可以使用 Nsight Compute 来比较内核函数的利用率。以下屏幕显示了这一点：

内核函数利用率的比较

如果我们希望追踪我们的优化工作并确定有效的组件，这将非常有用。

源视图

Nsight Compute 提供了各种我们可以调查的页面。其中一个有用的页面是 Source 页面。如果 CUDA 应用程序是使用-lineinfo选项构建的，Nsight Compute 可以显示与 CUDA C/C++源代码相关的信息和 CUDA SASS 代码。然后，我们可以分析瓶颈代码并调查它与 SASS 代码级别的关系。此外，它提供了一个 Live Registers 数字，以便我们可以调查内核函数中所需寄存器的数量。以下截图显示了 Source 页面：

如果您需要了解更多关于此功能的信息，您可以在此文档中找到相关信息-docs.nvidia.com/nsight-compute/NsightCompute/index.html#profiler-report-source-page。

Nsight Compute 提供了一个以 CUDA 内核性能分析为中心的操作，我们可以用来验证 Night Systems 和 Nsight Compute 具有不同的优化范围。

总结

在本章中，我们已经介绍了如何配置 GPU 应用程序并对其进行调试。了解这些 CUDA 工具将有助于您高效和有效地开发，因为它们可以帮助您找到瓶颈，并在短时间内找到错误和漏洞。

到目前为止，我们一直专注于单个 GPU 应用程序开发。然而，许多 GPU 应用程序使用多个 GPU 来实现更好的性能。在下一章中，我们将介绍如何编写能在多个 GPU 上运行并且具有可扩展性性能的代码。您将学习什么因素会影响性能以及如何实现良好的性能水平。您还将能够应用本章涵盖的工具来加强多 GPU 系统及其经验，解决下一章的问题。

第六章：可扩展的多 GPU 编程

到目前为止，我们一直致力于在单个 GPU 上获得最佳性能。密集节点与多个 GPU 已成为即将到来的超级计算机的迫切需求，特别是自从 ExaFLOP（每秒千亿次操作）系统成为现实以来。 GPU 架构具有高能效，因此近年来，具有 GPU 的系统在 Green500 榜单（www.top500.org/green500）中占据了大多数前十名。在 2018 年 11 月的 Green500 榜单中，前十名中有七个基于 NVIDIA GPU。

NVIDIA 的 DGX 系统现在在一个服务器中有 16 个 V100 32GB。借助统一内存和诸如 NVLink 和 NvSwitch 之类的互连技术，开发人员可以将所有 GPU 视为一个具有 512GB 内存的大型 GPU（16 个 GPU *每个 32GB）。在本章中，我们将深入讨论编写 CUDA 代码的细节，并利用 CUDA-aware 库在多 GPU 环境中实现节点内和节点间的可伸缩性。

在本章中，我们将涵盖以下主题：

• 使用高斯消元法解线性方程

• GPUDirect 点对点

• MPI 简介

• GPUDirect RDMA

• CUDA 流

• 额外的技巧

技术要求

本章需要一台带有现代 NVIDIA GPU（Pascal 架构或更高版本）的 Linux PC，并安装了所有必要的 GPU 驱动程序和 CUDA Toolkit（10.0 或更高版本）。如果您不确定您的 GPU 架构，请访问 NVIDIA GPU 网站（developer.nvidia.com/cuda-gpus）并确认您的 GPU 架构。本章的代码也可以在 GitHub 上找到：github.com/PacktPublishing/Learn-CUDA-Programming。

本章中的示例代码是使用 CUDA 版本 10.1 开发和测试的。但是，建议您使用最新版本（CUDA）或更高版本。

由于本章需要展示多 GPU 的交互，我们需要至少两个相同类型和架构的 GPU。还要注意，一些功能，如 GPUDirect RDMA 和 NVLink，仅支持 NVIDIA 的 Tesla 卡。如果您没有像 Tesla P100 或 Tesla V100 这样的 Tesla 卡，不要灰心。您可以安全地忽略其中一些功能。与我们在这里展示的情况相比，性能数字将会有所变化，但相同的代码将仍然有效。

在下一节中，我们将看一个示例，使用流行的高斯算法解决一系列线性方程，以演示如何编写多 GPU。

使用高斯消元法解线性方程

为了演示在节点内和节点间使用多个 GPU，我们将从一些顺序代码开始，然后将其转换为节点内和节点间的多个 GPU。我们将解决一个包含M个方程和N个未知数的线性方程组。该方程可以表示如下：

A × x = b

在这里，A是一个具有M行和N列的矩阵，x是一个列向量（也称为解向量），具有N行，b也是一个具有M行的列向量。找到解向量涉及在给定A和b时计算向量x。解线性方程组的标准方法之一是高斯消元法。在高斯消元法中，首先通过执行初等行变换将矩阵A减少为上三角矩阵或下三角矩阵。然后，通过使用回代步骤解决得到的三角形方程组。

以下伪代码解释了解线性方程所涉及的步骤：

1\. For iteration 1 to N (N: number of unknowns) 
    1.1 Find a row with non-zero pivot
    1.2 Extract the pivot row
    1.3 Reduce other rows using pivot row
2 Computing the solution vector through back substitution

让我们看一个示例，以便理解算法。假设方程组如下：

首先，我们将尝试设置基线系统，如下所示：

1. 准备您的 GPU 应用程序。此代码可以在本书的 GitHub 存储库中的06_multigpu/gaussian文件夹中找到。

2. 使用nvcc编译器编译您的应用程序，如下所示：

$ nvcc -o gaussian_sequential.out gaussian_sequential.cu
$ nvcc -o gaussian_single_gpu.out gaussian_single_gpu.cu
$ $ time ./gaussian_sequential.out
$ time ./gaussian_single_gpu.out

前面的步骤编译并运行了本章中存在的两个版本的代码：

• 顺序运行的 CPU 代码

• 在单个 GPU 上运行的 CUDA 代码

现在，让我们看看高斯消元的单 GPU 实现中的热点。

高斯消元的单 GPU 热点分析

让我们尝试理解和分析顺序和单 GPU 代码以建立基线。在此基础上，我们将增强并添加对多 GPU 运行的支持。

顺序 CPU 代码：以下代码显示了顺序实现的提取代码：

for( int n = 0; n < N; n++ ){
// M: number of equations, N: number of unknowns
    for( int pr = 0; pr < M; pr++ ){
        // finding the pivot row 
        //if pr satisfies condition for pivot i.e. is non zero 
        break; 
    }
    for( int r = 0; r < M; r++ ){
        // reduce all other eligible rows using the pivot row
        double ratio = AB[r*N+n]/AB[pr*N+n]
        for( int nn = n; nn < N + 1; nn++ ){
            AB[r * N + nn] -= (ratio*AB[pr * N + nn]);
        }
    }
}

从视觉上看，发生的操作如下：

在这里，高斯消元中的行数等于方程的数量，列数等于未知数的数量。在前面的图表中显示的pr行是主元行，将用于使用主元素减少其他行。

我们可以做出的第一个观察是，我们正在对增广矩阵进行操作，将A矩阵与b向量合并。因此，未知数的大小为N+1，因为增广矩阵的最后一列是b向量。创建增广矩阵有助于我们只处理一个数据结构，即矩阵。您可以使用以下命令对此代码进行分析。分析结果将显示guassian_elimination_cpu()函数完成所需的时间最长：

$ nvprof --cpu-profiling on ./guassian_sequential.out

CUDA 单 GPU 代码：通过前几章的学习，我们期望您已经熟悉了如何编写最佳的 GPU 代码，因此我们不会详细介绍单个 GPU 实现。以下摘录显示，在单个 GPU 实现中，三个步骤被称为三个用于找到N未知数的核心：

• findPivotRowAndMultipliers<<<...>>>：该核心查找主元行和乘数，应用于行消除。

• extractPivotRow<<<>>>：该核心提取主元行，然后用于执行行消除。

• rowElimination<<<>>>：这是最终的核心调用，在 GPU 上并行进行行消除。

以下代码片段显示了数据在复制到 GPU 后迭代调用的三个核心：

<Copy input augmented matrix AB to GPU>
...
for( int n = 0; n < N; n++ ){
// M: number of equations, N: number of unknowns
    findPivotRowAndMultipliers<<<...>>>(); 
    extractPivotRow<<<...>>>(); 
    rowElimination<<<...>>>(); 

}

本章的重点是如何增强此单个 GPU 实现以支持多个 GPU。但是，为了填补 GPU 实现中的缺失部分，我们需要对单个 GPU 实现进行一些优化更改：

• 高斯消元算法的性能受内存访问模式的影响很大。基本上，它取决于 AB 矩阵的存储方式：

• 找到主元行更喜欢列主格式，因为如果矩阵以列主格式存储，则提供了合并访问。

• 另一方面，提取主元行更喜欢行主格式。

• 无论我们如何存储AB矩阵，内存访问中都无法避免一个合并和一个跨步/非合并的访问。

• 列主格式对于行消除核心也是有益的，因此对于我们的高斯消元核心，我们决定存储 AB 矩阵的转置而不是 AB。AB 矩阵在代码开始时通过transposeMatrixAB()函数转置一次。

在下一节中，我们将启用多 GPU P2P 访问并将工作分配给多个 GPU。

GPU 直接点对点

GPUDirect 技术是为了允许 GPU 在节点内部和跨不同节点之间进行高带宽、低延迟的通信而创建的。该技术旨在消除一个 GPU 需要与另一个 GPU 通信时的 CPU 开销。GPUDirect 可以分为以下几个主要类别：

• GPU 之间的点对点（P2P）传输：允许 CUDA 程序在同一系统中的两个 GPU 之间使用高速直接内存传输（DMA）来复制数据。它还允许对同一系统中其他 GPU 的内存进行优化访问。

• 网络和存储之间的加速通信：这项技术有助于从第三方设备（如 InfiniBand 网络适配器或存储）直接访问 CUDA 内存。它消除了不必要的内存复制和 CPU 开销，从而减少了传输和访问的延迟。此功能从 CUDA 3.1 开始支持。

• 视频的 GPUDirect：这项技术优化了基于帧的视频设备的流水线。它允许与 OpenGL、DirectX 或 CUDA 进行低延迟通信，并且从 CUDA 4.2 开始支持。

• 远程直接内存访问（RDMA）：此功能允许集群中的 GPU 之间进行直接通信。此功能从 CUDA 5.0 及更高版本开始支持。

在本节中，我们将把我们的顺序代码转换为使用 GPUDirect 的 P2P 功能，以便在同一系统中的多个 GPU 上运行。

GPUDirect P2P 功能允许以下操作：

• GPUDirect 传输：cudaMemcpy()启动了从 GPU 1 的内存到 GPU 2 的内存的 DMA 复制。

• 直接访问：GPU 1 可以读取或写入 GPU 2 的内存（加载/存储）。

以下图表展示了这些功能：

要理解 P2P 的优势，有必要了解 PCIe 总线规范。这是为了通过 InfiniBand 等互连优化与其他节点进行通信而创建的。当我们想要从单个 GPU 优化地发送和接收数据时，情况就不同了。以下是一个样本 PCIe 拓扑，其中八个 GPU 连接到各种 CPU 和 NIC/InfiniBand 卡：

在前面的图表中，GPU0 和 GPU1 之间允许 P2P 传输，因为它们都位于同一个 PCIe 交换机中。然而，GPU0 和 GPU4 不能执行 P2P 传输，因为两个I/O Hub（IOHs）之间不支持 PCIe P2P 通信。IOH 不支持来自 PCI Express 的非连续字节进行远程对等 MMIO 事务。连接两个 CPU 的 QPI 链路的性质确保了如果 GPU 位于不同的 PCIe 域上，则不可能在 GPU 内存之间进行直接 P2P 复制。因此，从 GPU0 的内存到 GPU4 的内存的复制需要通过 PCIe 链路复制到连接到 CPU0 的内存，然后通过 QPI 链路传输到 CPU1，并再次通过 PCIe 传输到 GPU4。正如你所想象的那样，这个过程增加了大量的开销，无论是延迟还是带宽方面。

以下图表显示了另一个系统，其中 GPU 通过支持 P2P 传输的 NVLink 互连相互连接：

前面的图表显示了一个样本 NVLink 拓扑，形成了一个八立方网格，其中每个 GPU 与另一个 GPU 最多相连 1 跳。

更重要的问题是，*我们如何找出这个拓扑结构以及哪些 GPU 支持 P2P 传输？*幸运的是，有工具可以做到这一点。nvidia-smi就是其中之一，它作为 NVIDIA 驱动程序安装的一部分被安装。以下屏幕截图显示了在前面图表中显示的 NVIDIA DGX 服务器上运行nvidia-smi的输出：

前面的屏幕截图代表了在具有 8 个 GPU 的 DGX 系统上运行nvidia-smi topo -m命令的结果。如您所见，通过 SMP 互连（QPI/UPI）连接到另一个 GPU 的任何 GPU 都无法执行 P2P 传输。例如，GPU0将无法与GPU5、GPU6和GPU7进行 P2P 传输。另一种方法是通过 CUDA API 来找出这种传输，我们将在下一节中使用它来转换我们的代码。

现在我们已经了解了系统拓扑，我们可以开始将我们的应用程序转换为单个节点/服务器上的多个 GPU。

单节点-多 GPU 高斯消元

准备您的多 GPU 应用程序。此代码可以在本书的 GitHub 存储库中的06_multigpu/gaussian中找到。使用nvcc编译器编译您的应用程序，如下所示：

$ nvcc -o gaussian_multi_gpu_p2p.out gaussian_multi_gpu_p2p.cu
$ time ./gaussian_multi_gpu_p2p.out

从单 GPU 实现转换为多 GPU 实现，我们在上一小节中定义的三个内核将被原样使用。但是，线性系统被分成与 GPU 数量相等的部分。这些部分分配给每个 GPU 一个部分。每个 GPU 负责对分配给该 GPU 的部分执行操作。矩阵是按列分割的。这意味着每个 GPU 从所有行中获得相等数量的连续列。用于找到主元的内核在包含主元素的列上启动。主元元素的行索引被广播到其他 GPU。提取的主元行和行消除内核在所有 GPU 上启动，每个 GPU 都在矩阵的自己的部分上工作。以下图显示了行在多个 GPU 之间的分割以及主元行需要广播到其他进程的情况：

上述图表示了在多个 GPU 上的工作分配。目前，主元行属于GPU1，负责将主元行广播到其他 GPU。

让我们试着理解这些代码更改，以及用于启用 P2P 功能的 CUDA API：

1. 在支持的 GPU 之间启用 P2P 访问。以下代码显示了这个步骤的第一步：启用 GPU 之间的 P2P 访问：

for( int i = 0; i < nGPUs; i++ ){   
    // setup P2P 
    cudaSetDevice(i);   
    for( int j = 0; j < nGPUs; j++ ) {      
        if (i == j) continue;      
        cudaDeviceCanAccessPeer(&canAccessPeer, i, j);
        if (canAccessPeer)      
            cudaDeviceEnablePeerAccess(j, 0);    
    } 
}

在上述代码中使用的关键 API 如下：

• • cudaDeviceCanAccessPeer(): 检查当前 GPU 是否可以对传递的 GPU ID 进行 P2P 访问

• cudaDeviceEnablePeerAccess(): 如果cudaDeviceCanAccessPeer()返回True，则启用 P2P 访问

1. 拆分并将内容传输到各自的 GPU：

for( int g = 0; g < nGPUs; g++ ){       
    cudaSetDevice(g);       
    //Copy  part ‘g’ of ABT to GPU ‘g’; 
}

在上述代码中使用的关键 API 是cudaSetDevice()。这将当前上下文设置为作为参数传递的 GPU ID。

1. 找到主元行并通过 P2P 进行广播：

for( int n = 0; n < N; n++ ){        
    gp = GPU that holds n;        
    cudaSetDevice(gp);        
    findPivotRowAndMultipliers<<<...>>>();
    for( int g = 0; g < nGPUs; g++ ){ 
        if (g == gp) continue;
        cudaMemcpyPeer(pivotDatag, g, pivotDatagp, gp, numBytes);
     }  ... 

用于将传输广播到 GPU 的 API 是cudaMemcpyPeer()。

1. 提取主元行并执行行消除：

for( int n = 0; n < N; n++ ){
    ...
    for( int g = 0; g < nGPUs; g++ ){  
        cudaSetDevice(g); 
        extractPivotRow<<<...>>>(); 
        rowElimination<<<...>>>();   
    }  
}  

如您所见，我们仍在重用相同的内核。唯一的区别是我们使用cudaSetDevice() API 告诉 CUDA 运行时内核应该在哪个 GPU 上启动。请注意，cudaSetDevice()是一个昂贵的调用，特别是在旧一代的 GPU 上。因此，建议您通过在 CPU 上并行调用nGPUs的 for 循环，利用OpenMP/OpenACC或 CPU 上的任何其他线程机制来调用。

1. 从各自的 CPU 中复制数据回来：

for( int g = 0; g < nGPUs; g++ ){ 
    cudaSetDevice(g);  
    Copy  part ‘g’ of reduced ABT from GPU ‘g’ to Host; 
}

这五个步骤完成了将单个 GPU 实现转换为单个节点上的多个 GPU 的练习。

作为 CUDA 安装的一部分提供的 CUDA 示例包括一些测试 P2P 带宽性能的示例代码。它可以在samples/1_Utilities/p2pBandwidthLatencyTest文件夹中找到。建议您在系统上运行此应用程序，以便了解系统的 P2P 带宽和延迟。

现在我们已经在单个节点上实现了多 GPU，我们将改变方向并在多个 GPU 上运行此代码。但在将我们的代码转换为多个 GPU 之前，我们将提供一个关于 MPI 编程的简短介绍，这主要用于节点间通信。

MPI 的简要介绍

消息传递接口（MPI）标准是一种消息传递库标准，已成为在 HPC 平台上编写消息传递程序的行业标准。基本上，MPI 用于在多个 MPI 进程之间进行消息传递。相互通信的 MPI 进程可以驻留在同一节点上，也可以跨多个节点。

以下是一个 Hello World MPI 程序的示例：

#include <mpi.h> 
int main(int argc, char *argv[]) {     
    int rank,size;     
    /* Initialize the MPI library */     
    MPI_Init(&argc,&argv);     
    /* Determine the calling process rank and total number of ranks */
    MPI_Comm_rank(MPI_COMM_WORLD,&rank);     
    MPI_Comm_size(MPI_COMM_WORLD,&size);     
    /* Compute based on process rank */     
    /* Call MPI routines like MPI_Send, MPI_Recv, ... */     
    ...     
    /* Shutdown MPI library */     
    MPI_Finalize();     
    return 0; 
}

正如您所看到的，MPI 程序涉及的一般步骤如下：

1. 我们包括头文件mpi.h，其中包括所有 MPI API 调用的声明。

2. 我们通过调用MPI_Init并将可执行参数传递给它来初始化 MPI 环境。在这个语句之后，多个 MPI 等级被创建并开始并行执行。

3. 所有 MPI 进程并行工作，并使用诸如MPI_Send()、MPI_Recv()等消息传递 API 进行通信。

4. 最后，我们通过调用MPI_Finalize()终止 MPI 环境。

我们可以使用不同的 MPI 实现库（如 OpenMPI、MVPICH、Intel MPI 等）来编译此代码：

$ mpicc -o helloWorldMPI helloWorldMPI.c
$ mpirun -n 4 --hostfile hostsList ./helloWorldMPI

我们使用mpicc编译器来编译我们的代码。mpicc基本上是一个包装脚本，它在内部扩展编译指令，以包括相关库和头文件的路径。此外，运行 MPI 可执行文件需要将其作为参数传递给mpirun。mpirun是一个包装器，它帮助在应用程序应该执行的多个节点上设置环境。-n 4参数表示我们要运行四个进程，并且这些进程将在主机名存储在文件主机列表中的节点上运行。

在本章中，我们的目标是将 GPU 内核与 MPI 集成，使其在多个 MPI 进程中运行。但我们不会涵盖 MPI 编程的细节。那些不熟悉 MPI 编程的人应该先查看computing.llnl.gov/tutorials/mpi/，了解分布式并行编程，然后再进入下一节。

GPUDirect RDMA

在集群环境中，我们希望在多个节点上利用 GPU。我们将允许我们的并行求解器将 CUDA 代码与 MPI 集成，以利用多节点、多 GPU 系统上的多级并行性。使用 CUDA-aware MPI 来利用 GPUDirect RDMA 进行优化的节点间通信。

GPUDirect RDMA 允许在集群中的 GPU 之间进行直接通信。它首先由 CUDA 5.0 与 Kepler GPU 卡支持。在下图中，我们可以看到 GPUDirect RDMA，即Server 1中的GPU 2直接与Server 2中的GPU 1通信：

GPUDirect RDMA 工作的唯一理论要求是网络卡和GPU共享相同的根复杂性。 GPU 和网络适配器之间的路径决定了是否支持 RDMA。让我们重新访问我们在上一节中运行的 DGX 系统上nvidia-smi topo -m命令的输出：

如果我们看一下GPU4行，它显示GPU4到mlx5_2连接类型为PIX（通过 PCIe 交换机遍历）。我们还可以看到GPU4到mlx_5_0连接类型为SYS（通过QPI遍历）。这意味着GPU4可以通过 Mellanox InfiniBand 适配器mlx_5_2执行 RDMA 传输，但如果需要从mlx_5_0进行传输，则无法进行 RDMA 协议，因为QPI不允许。

CUDA-aware MPI

所有最新版本的 MPI 库都支持 GPUDirect 功能。支持 NVIDIA GPUDirect 和统一虚拟寻址（UVA）的 MPI 库使以下功能可用：

• MPI 可以将 API 传输直接复制到/从 GPU 内存（RDMA）。

• MPI 库还可以区分设备内存和主机内存，无需用户提示，因此对 MPI 程序员透明。

• 程序员的生产率提高了，因为少量应用代码需要更改以在多个 MPI 秩之间传输数据。

正如我们之前提到的，CPU 内存和 GPU 内存是不同的。没有 CUDA-aware MPI，开发人员只能将指向 CPU/主机内存的指针传递给 MPI 调用。以下代码是使用非 CUDA-aware MPI 调用的示例：

 //MPI rank 0:Passing s_buf residing in GPU memory 
 // requires it to be transferred to CPU memory
cudaMemcpy(s_buf_h,s_buf_d,size,cudaMemcpyDeviceToHost);
MPI_Send(s_buf_h,size,MPI_CHAR,1,100,MPI_COMM_WORLD);

//MPI rank 1: r_buf received buffer needs to be 
// transferred to GPU memory before being used in GPU
MPI_Recv(r_buf_h,size,MPI_CHAR,0,100,MPI_COMM_WORLD, &status);
cudaMemcpy(r_buf_d,r_buf_h,size,cudaMemcpyHostToDevice);

有了 CUDA-aware MPI 库，这是不必要的；GPU 缓冲区可以直接传递给 MPI，如下所示：

//MPI rank 0
MPI_Send(s_buf_d,size,MPI_CHAR,1,100,MPI_COMM_WORLD);

//MPI rank n-1
MPI_Recv(r_buf_d,size,MPI_CHAR,0,100,MPI_COMM_WORLD, &status);

例如，对于 Open MPI，CUDA-aware 支持存在于 Open MPI 1.7 系列及更高版本中。要启用此功能，需要在编译时配置 Open MPI 库以支持 CUDA，如下所示：

$ ./configure --with-cuda

拥有 CUDA-aware MPI 并不意味着总是使用 GPUDirect RDMA。如果数据传输发生在网络卡和 GPU 之间共享相同的根复杂，则使用 GPUDirect 功能。尽管如此，即使未启用 RDMA 支持，拥有 CUDA-aware MPI 也可以通过利用诸如消息传输之类的功能使应用程序更有效，如下图所示可以进行流水线处理：

！

上图显示了具有 GPUDirect 的 CUDA-aware MPI 与不具有 GPUDirect 的 CUDA-aware MPI。两个调用都来自 CUDA-aware MPI，但左侧是 GPUDirect 传输，右侧是没有 GPUDirect 传输。

非 GPUDirect 传输有以下阶段：

• 节点 1：从 GPU1 传输到主机内存

• 节点 1：从主机内存传输到网络适配器暂存区

• 网络：通过网络传输

• 节点 2：从网络暂存区传输到主机内存

• 节点 2：从主机内存传输到 GPU 内存

如果支持 GPUDirect RDMA，则从 GPU 传输直接通过网络进行，涉及主机内存的额外副本都被删除。

现在我们已经掌握了这个概念，让我们开始将代码转换为使用 CUDA-aware MPI 编程启用多 GPU 支持。

多节点-多 GPU 高斯消元

准备您的 GPU 应用程序。此代码可以在本书的 GitHub 存储库中的06_multigpu/gaussian中找到。使用nvcc编译器编译和运行应用程序，如下所示：

$ mpicc-o gaussian_multi_gpu_rdma.out gaussian_multi_gpu_rdma.cu
$ mpirun -np 8 ./gaussian_multi_gpu_rdma.out

我们使用mpicc而不是nvcc来编译 MPI 程序。我们使用mpirun命令运行可执行文件，而不是直接运行已编译的可执行文件。本节中您将看到的结果是在同一系统上具有 8 个 V100 的 DGX 系统上运行的输出。我们利用 8 个最大 MPI 进程，将每个 GPU 映射为 1 个 MPI 进程。要了解如何将多个 MPI 进程映射到同一 GPU，请阅读本章后面的MPS子节。在本练习中，我们使用了已编译为支持 CUDA 的 Open MPI 1.10，如前一节所述。

多 GPU 实现涉及的步骤如下：

1. MPI 进程的秩 0 生成线性系统（矩阵 A，B）的数据。

2. 转置增广矩阵（AB^T）由根节点在 MPI 进程之间使用MPI_Scatterv()按行分割。

3. 每个 MPI 进程并行计算其部分输入：

• 三个内核的处理发生在 GPU 上。

• 在findPivot操作后，通过MPI_Send()/Recv()实现了枢轴的共识。

1. 减少的转置增广矩阵（ABT）使用MPI_Gatherv()在根节点上收集。

2. 根节点执行回代以计算解 X。

展示前面代码的提取样本高斯代码如下：

void gaussianEliminationOnGPU() {
    cudaSetDevice(nodeLocalRank); //Set CUDA Device based on local rank
    //Copy  chuck of AB Transpose from Host to GPU; 
   for( int n = 0; n < N; n++ ){ 
       prank = MPI rank that holds n; 
       if (myRank == prank) 
           findPivotRowAndMultipliers<<<...>>>(); 
       bCastPivotInfo(); // from prank to other ranks 
       extractPivotRow<<<...>>>(); 
       rowElimination<<<...>>>(); 
   //Copy  myPartOfReducedTransposeAB from GPU to Host;
}

现在，让我们添加多 GPU 支持：

1. 设置每个 MPI 等级的 CUDA 设备：在 Open MPI 中，您可以通过使用MPI_COMM_TYPE_SHARED作为MPI_Comm_split_type的参数来获得 MPI 进程的本地等级，如下面的代码所示：

MPI_Comm loc_comm;
MPI_Comm_split_type(MPI_COMM_WORLD, MPI_COMM_TYPE_SHARED, rank, MPI_INFO_NULL, &loc_comm);
int local_rank = -1;
MPI_Comm_rank(loc_comm,&local_rank);
MPI_Comm_free(&loc_comm);

现在我们有了本地等级，每个 MPI 进程都使用它来通过cudaSetDevice()设置当前 GPU，如下图所示：

1. 使用MPI_Scatter将输入拆分并分发到不同的 MPI 进程：

void distributeInputs() {
    MPI_Scatterv(transposeAB, ..., myPartOfTransposeAB, recvCount, MPI_UNSIGNED, 0, MPI_COMM_WORLD); 
} 

1. 在 GPU 上执行高斯消元：

void gaussianEliminationOnGPU() { 
    cudaSetDevice(nodeLocalRank);
     for( int n = 0; n < N; n++ ){ 
        prank = MPI rank that holds n; 
        if (myRank == prank) 
            findPivotRowAndMultipliers<<<...>>>();
        MPI_Bcast(...); // from prank to other ranks 
        extractPivotRow<<<...>>>(); 
        rowElimination<<<...>>>(); 
}

在执行任何操作之前，基于本地等级设置当前 GPU。然后，由负责该行的进程提取枢轴行，然后将枢轴行广播到所有其他 MPI 等级，我们用于消除。

通过使用异步 MPI 调用而不是使用广播 API（如MPI_Bcast），可以提高传输时间的整体性能。实际上，不建议使用广播 API；它应该被替换为可以实现相同功能的MPI_Isend和MPI_Irecv，这些是异步版本。请注意，使调用异步会增加其他方面（如调试）的复杂性。因此，用户需要编写额外的代码来发送和接收数据。

本章提供了在向现有 MPI 程序添加 GPU 支持时的最佳编码实践，并不应被视为 MPI 编程的最佳编程实践的专家指南。

CUDA 流

流以 FIFO 方式工作，其中操作的顺序按照它们发出的顺序执行。从主机代码发出的请求被放入先进先出队列中。队列由驱动程序异步读取和处理，并且设备驱动程序确保队列中的命令按顺序处理。例如，内存复制在内核启动之前结束，依此类推。

使用多个流的一般想法是，在不同流中触发的 CUDA 操作可能会并发运行。这可能导致多个内核重叠或内核执行中的内存复制重叠。

为了理解 CUDA 流，我们将看两个应用程序。第一个应用程序是一个简单的矢量加法代码，添加了流，以便它可以重叠数据传输和内核执行。第二个应用程序是一个图像合并应用程序，也将在第九章中使用，使用 OpenACC 进行 GPU 编程。

首先，根据以下步骤配置您的环境：

1. 准备您的 GPU 应用程序。例如，我们将合并两个图像。此代码可以在本书的 GitHub 存储库的06_multi-gpu/streams文件夹中找到。

2. 使用nvcc编译器编译您的应用程序如下：

$ nvcc --default-stream per-thread -o vector_addition -Xcompiler -fopenmp -lgomp vector_addition.cu
$ nvcc --default-stream per-thread -o merging_muli_gpu -Xcompiler -fopenmp -lgomp scrImagePgmPpmPackage.cu image_merging.cu
$ ./vector addition
$ ./merging_muli_gpu

上述命令将创建两个名为vector_addition和merging_multi_gpu的二进制文件。正如您可能已经注意到的，我们在我们的代码中使用了额外的参数。让我们更详细地了解它们：

• --default-stream per-thread：此标志告诉编译器解析代码中提供的 OpenACC 指令。

• -Xcompiler -fopenmp -lgomp：此标志告诉nvcc将这些附加标志传递给 CPU 编译器，以编译代码的 CPU 部分。在这种情况下，我们要求编译器向我们的应用程序添加与 OpenMP 相关的库。

我们将把这一部分分为两部分。应用程序 1 和应用程序 2 分别演示了在单个和多个 GPU 中使用流。

应用程序 1-使用多个流来重叠数据传输和内核执行

我们需要遵循的步骤来重叠数据传输和内核执行，或者同时启动多个内核如下：

1. 声明要固定的主机内存，如下面的代码片段所示：

cudaMallocHost(&hostInput1, inputLength*sizeof(float));
cudaMallocHost(&hostInput2, inputLength*sizeof(float));
cudaMallocHost(&hostOutput, inputLength*sizeof(float));

在这里，我们使用cudaMallocHost() API 来分配固定内存的向量。

1. 创建一个Stream对象，如下面的代码片段所示：

for (i = 0; i < 4; i++) {
 cudaStreamCreateWithFlags(&stream[i],cudaStreamNonBlocking);

在这里，我们使用cudaStreamCreateWithFlags() API，传递cudaStreamNonBlocking作为标志，使此流非阻塞。

1. 调用 CUDA 内核和内存复制时使用stream标志，如下面的代码片段所示：

for (i = 0; i < inputLength; i += Seglen * 4) {
    for (k = 0; k < 4; k++) {
        cudaMemcpyAsync(... , cudaMemcpyHostToDevice, stream[k]);
        cudaMemcpyAsync(... , cudaMemcpyHostToDevice, stream[k]);
        vecAdd<<<Gridlen, 256, 0, stream[k]>>>(...);
    }
}

如我们所见，我们不是通过一次复制整个数组来执行矢量加法，而是将数组分成段，并异步复制这些段。内核执行也是在各自的流中异步进行的。

当我们通过 Visual Profiler 运行这段代码时，我们可以看到以下特点：

前面的分析器截图显示，蓝色条（基本上是vector_addition内核）重叠了内存复制。由于我们在代码中创建了四个流，分析器中也有四个流。

每个 GPU 都有两个内存复制引擎。一个负责主机到设备的传输，另一个负责设备到主机的传输。因此，发生在相反方向的两个内存复制可以重叠。此外，内存复制可以与计算内核重叠。这可以导致n路并发，如下图所示：

每个 GPU 架构都有一定的约束和规则，根据这些规则，我们将在执行时看到这些重叠。一般来说，以下是一些指导方针：

• CUDA 操作必须在不同的非 0 流中。

• 使用cudaMemcpyAsync时，主机应该使用cudaMallocHost()或cudaHostAlloc()进行固定。

• 必须有足够的资源可用。

• 不同方向的cudaMemcpyAsyncs

• 设备资源（SMEM、寄存器、块等）以启动多个并发内核

应用程序 2 - 使用多个流在多个设备上运行内核

为了在多个设备上运行内核并重叠内存传输，我们之前遵循的步骤保持不变，除了一个额外的步骤：设置 CUDA 设备以创建流。让我们看看以下步骤：

1. 创建与系统中 CUDA 设备数量相等的流，如下面的代码片段所示：

cudaGetDeviceCount(&noDevices);
cudaStream_t *streams;
streams = (cudaStream_t*) malloc(sizeof(cudaStream_t) * noDevices);

我们使用cudaGetDeviceCount() API 来获取 CUDA 设备的数量。

1. 在各自的设备中创建流，如下面的代码片段所示：

#pragma omp parallel num_threads(noDevices)
{
     int block = omp_get_thread_num();
    cudaSetDevice(block);
    cudaStreamCreate(&streams[block]);

我们启动与 CUDA 设备数量相等的 OpenMP 线程，以便每个 CPU 线程可以为其各自的设备创建自己的 CUDA 流。每个 CPU 线程执行cudaSetDevice()来根据其 ID 设置当前 GPU，然后为该设备创建流。

1. 在该流中启动内核和内存复制，如下所示：

cudaMemcpyAsync(... cudaMemcpyHostToDevice,streams[block]);
cudaMemcpyAsync(..., cudaMemcpyHostToDevice, streams[block]);
merging_kernel<<<gridDim,blockDim,0,streams[block]>>>(...);
cudaMemcpyAsync(...,streams[block]); 

在分析器中运行代码后的输出可以在下面的截图中看到，这代表了 Visual Profiler 的时间轴视图。这显示了一个 GPU 的内存复制与另一个 GPU 的内核执行重叠：

如您所见，我们在拥有四个 V100 的多 GPU 系统上运行了这段代码。不同 GPU 中的内存复制和内核重叠。在这段代码中，我们演示了利用 OpenMP 在不同设备上并行调用 CUDA 内核。这也可以通过利用 MPI 来启动利用不同 GPU 的多个进程来实现。

在下一节中，我们将看一些额外的主题，这些主题可以提高多 GPU 应用程序的性能，并帮助开发人员分析和调试他们的代码。

额外的技巧

在本节中，我们将涵盖一些额外的主题，这些主题将帮助我们了解多 GPU 系统的额外特性。

使用 InfiniBand 网络卡对现有系统进行基准测试

有不同的基准可用于测试 RDMA 功能。InfiniBand 适配器的一个这样的基准可以在www.openfabrics.org/找到。您可以通过执行以下代码来测试您的带宽：

$ git clone git://git.openfabrics.org/~grockah/perftest.git
$ cd perftest 
$ ./autogen.sh 
$ export CUDA_H_PATH=<<Path to cuda.h>> 
$ ./configure –prefix=$HOME/test 
$ make all install

然后，您可以运行以下命令来测试带宽：

For example host to GPU memory (H-G) BW test:
server$ ~/test/bin/ib_write_bw -n 1000 -O -a --use_cuda
client $ ~/test/bin/ib_write_bw -n 1000 -O -a server.name.org

//GPU to GPU memory (G-G) BW test:
server$ ~/test/bin/ib_write_bw -n 1000 -O -a --use_cuda
client $ ~/test/bin/ib_write_bw -n 1000 -O -a --use_cuda server.name.org

NVIDIA 集体通信库（NCCL）

NCCL 提供了常用于深度学习等领域的通信原语的实现。NCCL 1.0 从同一节点内多个 GPU 之间的通信原语实现开始，并发展到支持多个节点上的多个 GPU。NCCL 库的一些关键特性包括以下内容：

• 支持来自多个线程和多个进程的调用

• 支持多个环和树拓扑，以更好地利用节点内和节点间的总线

• 支持 InfiniBand 节点间通信

• 源代码包可以从 GitHub（github.com/nvidia/nccl）免费下载

NCCL 可以扩展到 24,000 个 GPU，延迟低于 300 微秒。请注意，尽管 NCCL 已被证明是深度学习框架中非常有用和方便的库，但在用于 HPC 应用时存在局限，因为它不支持点对点通信。NCCL 支持集体操作，这在深度学习应用中被使用，例如以下内容：

• AllReduce

• AllGather

• ReduceScatter

• Reduce

• Broadcast

所有 NCCL 调用都作为 CUDA 内核运行，以更快地访问 GPU 内存。它使用较少的线程，实现为一个块。这最终只在一个 GPU SM 上运行，因此不会影响其他 GPU 的利用率。让我们看一下以下代码：

ncclGroupStart(); 
for (int i=0; i<ngpus; i++) 
{ 
    ncclAllGather(…, comms[i], streams[i]); 
} 
ncclGroupEnd();

正如我们所看到的，NCCL 调用简单，易于调用。

使用 NCCL 加速集体通信

**NVIDIA 集体通信库（NCCL）**提供了为多个 NVIDIA GPU 优化的性能集体通信原语。在本节中，我们将看到这个库是如何工作的，以及我们如何从中受益。

并不难找到使用多个 GPU 来训练网络的深度学习模型。由于两个 GPU 并行计算神经网络，我们很容易想象这种技术将随着 GPU 数量的增加而提高训练性能。不幸的是，世界并不那么简单。梯度应该在多个 GPU 之间共享，并且一个 GPU 中的权重更新过程应该等待其他 GPU 的梯度来更新其权重。这是使用多个 GPU 进行深度学习训练的一般过程，并在以下图表中显示：

集体通信有许多类型：全局归约、广播、归约、全局收集、归约散射等。在深度学习中，每个 GPU 在传输自己的数据的同时收集另一个 GPU 的数据。因此，我们可以确定深度学习在通信中需要所有类型的归约样式通信。

在 HPC 社区中，包括全局归约在内的集体通信是一个常见的话题。节点内和节点间处理器之间的通信是一个具有挑战性但至关重要的问题，因为它直接关系到可扩展性。正如我们在第六章中提到的，可扩展的多 GPU 编程，在多 GPU 编程部分，需要仔细考虑与每个 GPU 的通信。开发人员应该设计和实现 GPU 中的集体通信，即使 MPI 已经支持这样的通信模式。

NCCL 提供了一种集体通信，它了解 GPU 拓扑配置。通过使用各种分组和通信命令，您可以应用所需的通信任务。

一个前提是您的系统需要有多个 GPU，因为 NCCL 是一个与多个 GPU 一起工作的通信库。

以下步骤涵盖了如何调用ncclAllReduce()来测试和测量系统的 GPU 网络带宽。示例代码实现在04_nccl中：

1. 让我们定义一个类型，它将包含、发送和接收每个 GPU 设备的缓冲区和cudaStream，如下所示：

typedef struct device
{
    float *d_send;
    float *d_recv;
    cudaStream_t stream;
} device_t;

1. 在应用程序开始时，我们需要准备一些句柄，以便我们可以控制多个 GPU：

cudaGetDeviceCount(&num_dev);
ncclComm_t *ls_comms = new ncclComm_t[num_dev];
int *dev_ids = new int[num_dev];
for (int i = 0; i < num_dev; i++)
    dev_ids[i] = i;

1. 然后，我们将创建一个缓冲区，假设我们有数据。对于每个设备，我们将初始化每个设备的项目，如下所示：

unsigned long long size = 512 * 1024 * 1024; // 2 GB

// allocate device buffers and initialize device handles
device_t *ls_dev = new device_t[num_dev];
for (int i = 0; i < num_dev; i++) {
    cudaSetDevice(i);
    cudaMalloc((void**)&ls_dev[i].d_send, sizeof(float) * size);
    cudaMalloc((void**)&ls_dev[i].d_recv, sizeof(float) * size);
    cudaMemset(ls_dev[i].d_send, 0, sizeof(float) * size);
    cudaMemset(ls_dev[i].d_recv, 0, sizeof(float) * size);
    cudaStreamCreate(&ls_dev[i].stream);
}

1. 在开始 NCCL 通信之前，我们需要初始化 GPU 设备，以便它们知道它们在 GPU 组中的排名。由于我们将用单个进程测试带宽，我们可以安全地调用一个初始化所有设备的函数：

ncclCommInitAll(ls_comms, num_dev, dev_ids);

1. 如果我们要用多个进程测试带宽，我们需要调用ncclCommInitRank()。我们需要为计算进程 ID 和 GPU 排名提供 GPU ID。

2. 现在，我们可以使用 NCCL 完成 all-reduce 操作。以下代码是ncclAllReduce的示例实现：

ncclGroupStart();
for (int i = 0; i < num_dev; i++) {
    ncclAllReduce((const void*)ls_dev[i].d_send, 
                  (void*)ls_dev[i].d_recv,
        test_size, ncclFloat, ncclSum, 
        ls_comms[i], ls_dev[i].stream);
}
ncclGroupEnd();

对于每个设备，我们需要触发流量。为此，我们需要启动和关闭 NCCL 组通信。现在，我们已经实现了一些使用ncclAllReduce()的测试代码。让我们通过微基准测试来了解 NCCL 的工作原理。

在多 GPU 系统上测试此代码，运行以下命令：

$ nvcc -run -m64 -std=c++11 -I/usr/local/cuda/samples/common/inc -gencode arch=compute_70,code=sm_70 -lnccl -o nccl ./nccl.cu

以下图表显示了在 DGX Station 中使用四个 V100 32G GPU 测得的性能。蓝线表示基于 NVLink 的带宽，而橙线表示基于 PCIe 的带宽，通过设置NCCL_P2P_DISABLE=1 ./ncd并关闭对等 GPU 来实现：

这个 NCCL 测试可能会受到系统配置的影响。这意味着结果可能会有所不同，取决于您系统的 GPU 拓扑结构。

这显示了基于 PCI Express 和基于 NVLINK 的 all-reduce 性能差异。我们可以使用nvprof来查看通信。以下屏幕截图显示了通过 NCCL 2.3.7 在 DGX Station 上的 all-reduce 通信：

NCCL 越来越快。通过引入新的 GPU 互连技术 NVLink 和 NVSwitch，我们对 NCCL 的经验正在增加，以至于我们可以实现可扩展的性能。

以下链接提供了关于 NCCL 的讨论：developer.nvidia.com/gtc/2019/video/S9656/video。

摘要

在本章中，我们介绍了多 GPU 编程的不同方法。通过示例高斯消元，我们看到了如何将单个 GPU 应用程序工作负载分割到多个 GPU 中，首先是单个节点，然后是多个节点。我们看到了系统拓扑在利用 P2P 传输和 GPUDirect RDMA 等功能方面起着重要作用。我们还看到了如何使用多个 CUDA 流来重叠多个 GPU 之间的通信和数据传输。我们还简要介绍了一些其他主题，可以帮助 CUDA 程序员优化代码，如 MPS 和使用nvprof来分析多 GPU 应用程序。

在下一章中，我们将看到大多数 HPC 应用程序中出现的常见模式以及如何在 GPU 中实现它们。