[NVIDIA]Accelerating Applications with CUDA C/C++笔记

void CPUFunction()
{
  printf("This function is defined to run on the CPU.\n");
}

__global__ void GPUFunction()
{
  printf("This function is defined to run on the GPU.\n");
}

int main()
{
  CPUFunction();

  GPUFunction<<<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 上恢复执行。
 

!nvcc -o hello-gpu 01-hello/01-hello-gpu.cu -run

执行配置允许程序员指定有关启动内核以在多个 GPU 线程上并行运行的详细信息。更准确地说，执行配置允许程序员指定线程组（称为线程块或块）的数量，以及他们希望每个线程块包含多少个线程。其语法为：

<<< NUMBER_OF_BLOCKS，NUMBER_OF_THREADS_PER_BLOCK>>>

** 内核代码由内核启动时配置的每个线程块中的每个线程执行**。

线程和块索引
每个线程在其线程块中都有一个索引，从 0 开始。此外，每个块都有一个索引，从 0 开始。正如线程被分组为线程块一样，块也被分组到网格中，网格是 CUDA 线程层次结构中的最高实体。总之，CUDA 内核在 1 个或多个块的网格中执行，每个块包含相同数量的 1 个或多个线程。

CUDA 内核可以访问特殊变量，这些变量既可以识别执行内核的线程（在块内）的索引，也可以标识线程所在的块（在网格内）的索引。这些变量分别是 threadIdx.x 和 blockIdx.x。
 

// CPU-only

int N = 2<<20;
size_t size = N * sizeof(int);

int *a;
a = (int *)malloc(size);

// Use `a` in CPU-only program.

free(a);
// Accelerated

int N = 2<<20;
size_t size = N * sizeof(int);

int *a;
// Note the address of `a` is passed as first argument.
cudaMallocManaged(&a, size);

// Use `a` on the CPU and/or on any GPU in the accelerated system.

cudaFree(a);

处理块配置与所需线程数量的不匹配
 

```cpp
// Assume `N` is known
int N = 100000;

// Assume we have a desire to set `threads_per_block` exactly to `256`
size_t threads_per_block = 256;

// Ensure there are at least `N` threads in the grid, but only 1 block's worth extra
size_t number_of_blocks = (N + threads_per_block - 1) / threads_per_block;

some_kernel<<<number_of_blocks, threads_per_block>>>(N);
```

```cpp
__global__ some_kernel(int N)
{
  int idx = threadIdx.x + blockIdx.x * blockDim.x;

  if (idx < N) // Check to make sure `idx` maps to some value within `N`
  {
    // Only do work if it does
  }
}
```

CUDA 提供了一个特殊变量，用于给出网格中的块数 gridDim.x。然后计算网格中的线程总数，只需将网格中的块数乘以每个块中的线程数，gridDim.x * blockDim.x。考虑到这一点，下面是一个内核中网格步幅循环的详细示例：
 

__global__ void kernel(int *a, int N)
{
  int indexWithinTheGrid = threadIdx.x + blockIdx.x * blockDim.x;
  int gridStride = gridDim.x * blockDim.x;

  for (int i = indexWithinTheGrid; i < N; i += gridStride)
  {
    // do work on a[i];
  }
}

Error Handling

cudaError_t err;
err = cudaMallocManaged(&a, N)                    // Assume the existence of `a` and `N`.

if (err != cudaSuccess)                           // `cudaSuccess` is provided by CUDA.
{
  printf("Error: %s\n", cudaGetErrorString(err)); // `cudaGetErrorString` is provided by CUDA.
}

/*
 * This launch should cause an error, but the kernel itself
 * cannot return it.
 */

someKernel<<<1, -1>>>();  // -1 is not a valid number of threads.

cudaError_t err;
err = cudaGetLastError(); // `cudaGetLastError` will return the error from above.
if (err != cudaSuccess)
{
  printf("Error: %s\n", cudaGetErrorString(err));
}

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

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

int main()
{

/*
 * The macro can be wrapped around any function returning
 * a value of type `cudaError_t`.
 */

  checkCuda( cudaDeviceSynchronize() )
}

dim3 threads_per_block(16, 16, 1);
dim3 number_of_blocks(16, 16, 1);
someKernel<<<number_of_blocks, threads_per_block>>>();