AC_CUDA_C-Copy1

练习：编写一个Hello GPU核函数¶

CUDA 为许多常用编程语言提供扩展，而在本实验中，我们将会为 C/C++ 提供扩展。这些语言扩展可让开发人员在 GPU 上轻松运行其源代码中的函数。

以下是一个 .cu 文件（.cu 是 CUDA 加速程序的文件扩展名）。其中包含两个函数，第一个函数将在 CPU 上运行，第二个将在 GPU 上运行。请抽点时间找出这两个函数在定义方式和调用方式上的差异。

void helloCPU()
{
  printf("Hello from the CPU.\n");
}

__global__ void helloGPU()
{
  printf("Hello also from the GPU.\n");
}

int main()
{
  helloCPU();
  helloGPU<<<1, 1>>>();
  cudaDeviceSynchronize();
}

以下是一些需要特别注意的重要代码行，以及加速计算中使用的一些其他常用术语：

__global__ void GPUFunction()

__global__ 关键字表明以下函数将在 GPU 上运行并可全局调用，而在此种情况下，则指由 CPU 或 GPU 调用。
通常，我们将在 CPU 上执行的代码称为主机代码，而将在 GPU 上运行的代码称为设备代码。
注意返回类型为 void。使用 __global__ 关键字定义的函数需要返回 void 类型。

GPUFunction<<<1, 1>>>();

通常，当调用要在 GPU 上运行的函数时，我们将此种函数称为已启动的核函数。
启动核函数时，我们必须提供执行配置，即在向核函数传递任何预期参数之前使用 <<< ... >>> 语法完成的配置。
在宏观层面，程序员可通过执行配置为核函数启动指定线程层次结构，从而定义线程组（称为线程块）的数量，以及要在每个线程块中执行的线程数量。稍后将在本实验深入探讨执行配置，但现在请注意正在使用包含 1 线程（第二个配置参数）的 1 线程块（第一个执行配置参数）启动核函数。

cudaDeviceSynchronize();

与许多 C/C++ 代码不同，核函数启动方式为异步：CPU 代码将继续执行而无需等待核函数完成启动。
调用 CUDA 运行时提供的函数 cudaDeviceSynchronize 将导致主机 (CPU) 代码暂作等待，直至设备 (GPU) 代码执行完成，才能在 CPU 上恢复执行。 ```

请进行以下修改，并尝试在每次更改后编译并运行该应用程序。若出现错误，请花时间仔细阅读错误内容：熟知错误内容会在您开始编写自己的加速代码时，给与很大的帮助。

从核函数定义中删除关键字 __global__。注意错误中的行号：您认为错误中的 \”configured\” 是什么意思？完成后，请替换 __global__。

>>>
01-hello/01-hello-gpu.cu(29): error: a host function call cannot be configured

1 error detected in the compilation of "/tmp/tmpxft_0000009f_00000000-8_01-hello-gpu.cpp1.ii".

启动核函数时，我们必须提供执行配置, 但是删除global后的函数实际上是普通的host function，无法进行configure。

移除执行配置：您对 \”configured\” 的理解是否仍旧合理？完成后，请替换执行配置。

>>>
01-hello/01-hello-gpu.cu(29): error: a __global__ function call must be configured

1 error detected in the compilation of "/tmp/tmpxft_000000a8_00000000-8_01-hello-gpu.cpp1.ii".

移除对 cudaDeviceSynchronize 的调用。在编译和运行代码之前，猜猜会发生什么情况，可以回顾一下核函数采取的是异步启动，且 cudaDeviceSynchronize 会使主机执行暂作等待，直至核函数执行完成后才会继续。完成后，请替换对 cudaDeviceSynchronize 的调用。

A: 程序顺利执行完毕，但未输出GPU的函数的输出。

重构 01-hello-gpu.cu，以便 Hello from the GPU 在 Hello from the CPU 之前打印。

#include <stdio.h>

void helloCPU()
{
  printf("Hello from the CPU.\n");
}

__global__ void helloGPU()
{
  printf("Hello also from the GPU.\n");
}

int main()
{

  helloGPU<<<1, 1>>>();

  cudaDeviceSynchronize();

  helloCPU();
}

重构 01-hello-gpu.cu，以便 Hello from the GPU 打印两次，一次是在 Hello from the CPU 之前，另一次是在 Hello from the CPU 之后。

#include <stdio.h>

void helloCPU()
{
  printf("Hello from the CPU.\n");
}

__global__ void helloGPU()
{
  printf("Hello also from the GPU.\n");
}

int main()
{

  helloGPU<<<1, 1>>>();

  cudaDeviceSynchronize();

  helloCPU();

  helloGPU<<<1, 1>>>();

  cudaDeviceSynchronize();
}

使用 CUDA C/C++ 加速应用程序

练习：编写一个Hello GPU核函数¶