《CUDA Programming Guide》2. CUDA C++

Basic Usage of CUDA C++

使用 __global__ 来标识一个 Kernel（前面说过入口 Device Function 才叫做 Kernel），__device__、__host__ 分别标识 Device Function 和 Host Function

使用三重尖括号 kernel_name<<<gridDim, blockDim>>>(...) 来启动 Kernel，其中两个 Dim 可以是数字也可以是 dim3（CUDA 内置类型），表示前面提到的 Grid-Block-Thread 的三级结构

在 Device 函数中使用 gridDim、blockDim 来表示 Grid 和 Block 的维度，blockIdx、threadIdx 表示 Block 和 Thread 在 Grid、Block 中的 Index，它们都有 .x、.y、.z 三个成员

Host 函数中，使用 cudaMalloc/cudaFree 来进行申请/释放显存，使用 cudaMallocHost/cudaFreeHost 来申请/释放 Page-Locked 内存（内存页表不会交换，显存内存之间传输数据时会更快），使用 cudaMemcpy 来进行数据传输（内存到显存或者显存到内存），使用 cudaMemset 来初始化显存

Host 函数中，也可以使用 cudaMallocManaged 来使用 CUDA 的 Unified Memory 特性，让驱动程序来管理 Host 与 Device 之间的数据移动（注意完成后需要使用 cudaFree 进行释放）

Host 函数中，使用 cudaDeviceSynchronize 来阻塞 Host 线程，直到 GPU 上所有先前提交的任务全部完成

Device 函数中，使用 __syncthreads() 来进行同一个 Block 内线程的同步

Device 函数中，变量也有特定的修饰符用来表示变量存放的位置，__device__、__constant__、__shared__ 分别表示变量会存放在 Global Memory、Constant Memory、Shared Memory 中，另外 __managed__ 表示变量是 Unified Memory 的，驱动程序来管理

示例代码：

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

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

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

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

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

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

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

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

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

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

  // Launch the kernel
  int threads = 256;
  int blocks = cuda::ceil_div(vectorLength, threads);
  vecAdd<<<blocks, threads>>>(devA, devB, devC, vectorLength);

  // wait for kernel execution to complete
  cudaDeviceSynchronize();

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

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

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

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


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

Writting SIMT Kernels

Memory Spaces

CUDA 设备具有多个可由 CUDA 线程访问的内存空间：

Memory Type	Scope	Lifetime	Location
Global	Grid	Application	Device
Constant	Grid	Application	Device
Shared	Block	Kernel	SM
Local	Thread	Kernel	Device
Register	Thread	Kernel	SM

Global Memory 是所有 Thread 都可访问的主要设备内存，一般通过 cudaMalloc 来分配，容量大、生命周期长，但是访问延迟通常较高

Shared Memory 由 Block 内所有线程共享，容量小但是延迟更低、带宽更高，可以动态或者静态分配

静态分配：Device Function 中，变量写 __shared__ 修饰符，比如：__shared__ float tile[32][32];
动态分配：Device Function 中，写 extern __shared__ shared[];，然后在启动 Kernel 时指定字节数，即 kernel<<<grid, block, sharedBytes>>>();
注意一个 Kernel 内所有 extern __shared__ 声明实际上引用同一块动态 Shared Memory，需要多个数组时，应手动划分并保证地址对齐
静态分配和动态分配可以同时存在

Local Memory 位于 Global Memory 地址空间，但它是 Thread 私有的，由编译器管理。实际上 Local Memory 充当了 CPU 编程中的“栈空间”功能，对于简单的局部变量，GPU 会放在寄存器中，而对于更复杂的变量（例如无法使用常量索引分析的局部数组、较大的局部结构体或数组、寄存器不足时溢出的变量）容易被编译器放进 Local Memory

Constant Memory 只能在 Kernel 或者 Device Function 之外声明。适用于存储少量数据，且这些数据需供各线程以只读方式使用。容量较小，通常每个设备仅有 64KB。Host 端可通过运行时库（如 cudaGetSymbolAddress()、cudaGetSymbolSize()、cudaMemcpyToSymbol()、cudaMemcpyFromSymbol()）进行访问

Memory Performance

Global Memory 只能一次性传输 32 字节，在这里每次传输被叫做 32-Byte Memory Transaction（内存事务）

在 Thread 请求数据时，Warp 会将所有的 Thread 的内存请求合并（Coalesce）为若干个内存事务，然后再进行数据传输

可以用内存利用率（线程真正使用的字节数除以实际传输的字节数）来描述内存合并访问的效率。理想情况下，所有 Thread 访问连续的地址，那内存利用率可达到 $100\%$；最差情况下，每个线程请求的地址相隔至少 32 字节，此时内存利用率就是 $1/32=12.5\%$

Coalescing 不等于“线程必须严格按顺序访问”，只要这些地址仍集中在相同的少数几个 32 字节块中，事务数量就不会增加

此外，对齐也会影响内存利用率

朴素的矩阵转置的代码可以写成：

__global__ void transposeNaive(float* output, const float* input, int n, int cols, int rows) {
  const auto col = blockIdx.x * blockDim.x + threadIdx.x; // threadIdx.x 是连续的，blockIdx 不变，所以 col 是连续的
  const auto row = blockIdx.y * blockDim.y + threadIdx.y; // threadIdx.y 不变，所以 row 是不变的

  if (row < rows && col < cols) {
    output[row + rows * col] = input[col + cols * row]; // 前面的不连续，后面的连续
  }
}

这里读取是合并的，但是写入不是合并的

可以对矩阵分块，然后进行转置，利用 Shared Memory：

先对矩阵 $\boldsymbol{A}$ 分块，每个子矩阵 $\boldsymbol{A}_{ij}$ 对应一个 Block，这是一个二维的
每个子矩阵 $\boldsymbol A_{ij}$ 中的一个元素对应一个 Thread，这也是二维的
$\begin{bmatrix}\boldsymbol{A}_{ij}\end{bmatrix}^T=\begin{bmatrix}\boldsymbol{A}_{ji}^T\end{bmatrix}$

设置子矩阵的大小为 $32\times 32$，这样 blockDim.x = blockDim.y = 32，先把读出来的东西写入 Shared Memory，转置后再进行连续的写入

__global__ void transposeNaive(float* output, const float* input, int cols, int rows) {
  __shared__ float tile[32][32]; // 行主序，即 tile[i][j] == tile[i * 32 + j]

  const auto input_col = blockIdx.x * 32 + threadIdx.x;
  const auto input_row = blockIdx.y * 32 + threadIdx.y;
    
  if (input_col < cols && input_row < rows) {
    tile[threadIdx.y][threadIdx.x] = input[input_row * cols + input_col]; // 写入 Shared Memory 然后转置
  }

  __syncthreads(); // 同步，等子矩阵中的元素全部读入并转置完成

  const auto output_col = blockIdx.y * 32 + threadIdx.x;
  const auto output_row = blockIdx.x * 32 + threadIdx.y;

  if (output_row < cols && output_col < rows) {
    output[output_row * rows + output_col] = tile[threadIdx.x][threadIdx.y]; // 写入输出矩阵，这是连续的
  }
}

可以发现，这样的话，转置操作被放在了 Shared Memory 中，从原矩阵到 Shared Memory 和从 Shared Memory 到目标矩阵都是复制操作，都是访存合并的

但是，这样会有 Bank Conflict

为了让一个 Warp 的 32 个线程可以并行访问 Shared Memory，需要把 Shared Memory 分成若干个存储端口，即 32 个 Bank：

Shared Memory
-------------------------------------------------------
| Bank 0 | Bank 1 | Bank 2 | Bank 3 |  ...  | Bank 31 |
-------------------------------------------------------

每个 Bank 可以独立处理访问。这样，一个 Warp 的 32 个线程如果分别访问不同 Bank，就可以同时完成

理想情况下，Warp 中 32 个线程分别访问 32 个不同 Bank

在上面的代码中有两处存取 Shared Memory：

tile[threadIdx.x][threadIdx.y]：相当于 threadIdx.x * 32 + threadIdx.y，其中 threadIdx.x 连续，threadIdx.y 不变，对 32 取余后不变，这样一个 Warp 中的所有线程访问了同一个 Bank（特殊情况是都访问同一个位置，此时是 Broadcast 机制，没有 Bank Conflict）
tile[threadIdx.y][threadIdx.x]：相当于 threadIdx.y * 32 + threadIdx.x，一个 Warp 中所有线程都访问的不同 Bank

如果 Warp 中的所有线程访问了同一个 Bank，那么访存操作全部是串行的，这就是 Bank Conflict

一个最简单的解决 Bank Conflict 的方法就是加一个 Padding，即 __shared__ float tile[32][32 + 1];，这样 tile[threadIdx.x][threadIdx.y] 就变成了 threadIdx.x * 32 + threadIdx.y + threadIdx.x，对 32 取余后各个线程就不同了

在《CUDA Programming Guide》中，为了减小索引计算成本，用一个线程处理 4 个元素

代码大概长这个样子：

__global__ void transposeNaive(float* output, const float* input, int cols, int rows) {
  __shared__ float tile[32][32 + 1];

  const auto input_col = blockIdx.x * 32 + threadIdx.x;
  const auto input_row = blockIdx.y * 32 + threadIdx.y;

  if (input_col < cols) {
    #pragma unroll
    for (int i = 0; i < 32; i += 8) {
      const auto row = input_row + i;
      if (row >= rows) break;
      tile[threadIdx.y + i][threadIdx.x] = input[row * cols + input_col];
    }
  }

  __syncthreads();

  const auto output_col = blockIdx.y * 32 + threadIdx.x;
  const auto output_row = blockIdx.x * 32 + threadIdx.y;

  if (output_col < rows) {
    #pragma unroll
    for (int i = 0; i < 32; i += 8) {
      const auto row = output_row + i;
      if (row >= cols) break;
      output[row * rows + output_col] = tile[threadIdx.x][threadIdx.y + i];
    }
  }
}

然后 blockDim.x = 32、blockDim.y = 8

Atomics

CUDA C++ 提供多种 Atomic 接口，比如类似 C++ 的接口：

cuda::atomic<int, cuda::thread_scope_block>：本身管理内存，然后第二个参数可以指定 CUDA Thread Scope，
cuda::atomic_ref<int, cuda::thread_scope_device>：不拥有数据，而是把一个已经存在的普通对象包装成 Atomic 引用

常见 Scope 有：

cuda::thread_scope_block：只需要在同一个 Block 的线程之间保持原子性
cuda::thread_scope_device：在同一个 GPU 的线程之间保持原子性
cuda::thread_scope_system：在所有 CPU 和 GPU 之间保持原子性，但是需要设备支持

Atomic 可以支持 Global Memory 和 Shared Memory

Asynchronous Execution

CUDA 允许 CPU、GPU 上的计算、内存与显存间的数据传输并行执行，这种并发性通过异步接口实现

通常，异步接口提供三种主要方式来与已经 dispatched 的操作进行同步：

Blocking Approach：比如 cudaDeviceSynchronize()，这个函数会阻塞等待前面所有已经开始的 GPU 计算完成，还有 cudaMemcpy(...) 也是阻塞式的
Non-blocking Approach or Polling Approach：开始任务立刻返回一个状态函数，比如 Kernal Launching 都是非阻塞式的，还有后面提到的 Stream 相关的函数
Callback Approach：操作完成后执行预先注册的函数

CUDA Stream

类似任务队列，先进先处理，这样使用：

// Allocate memory

cudaStream_t stream;        // Stream handle
cudaStreamCreate(&stream);  // Create a new stream

cudaMemcpyAsync(d, h, bytes, cudaMemcpyHostToDevice, stream);     // Launching Memory Transfers in CUDA Streams
kernel<<<grid, block, shared_mem_size, stream>>>(...);            // Launching Kernels in CUDA Streams
cudaMemcpyAsync(h_out, d, bytes, cudaMemcpyDeviceToHost, stream);

// Have a peek at the stream
// returns cudaSuccess if the stream is empty
// returns cudaErrorNotReady if the stream is not empty
cudaError_t status = cudaStreamQuery(stream);
switch (status) {
  case cudaSuccess:
    // The stream is empty
    std::cout << "The stream is empty" << std::endl;
    break;
  case cudaErrorNotReady:
    // The stream is not empty
    std::cout << "The stream is not empty" << std::endl;
    break;
  default:
      // An error occurred - we should handle this
    break;
};


// Wait for the stream to be empty of tasks
cudaStreamSynchronize(stream);

// At this point the stream is done
// and we can access the results of stream operations safely
...

cudaStreamDestroy(stream);     // Destroy the stream

上面的代码中：

cudaMemcpyAsync(...) 函数比 cudaMemcpy(...) 多了一个 stream 参数，加入到流中。其他的内存传输函数也有相应的异步版本
启动 Kernel 时，可以在三尖括号中传入 stream 参数，把这个 Kernel 加入到流中
cudaStreamQuery(stream) 函数是轮询方法，查看流的状态
cudaStreamSynchronize(stream) 阻塞并等待流中所有工作全部完成

CUDA 还提供了 Stream 的 Callback Function 支持，使用 cudaLaunchHostFunc() 来使用：

void host_function(void* data) {
  ...
}

// Sig: cudaError_t cudaLaunchHostFunc(cudaStream_t stream, void (*func)(void *), void *data);
cudaLaunchHostFunc(stream, host_function, user_data);

注意，Callback Function 中不能调用 CUDA API，否则可能产生不合法行为或死锁风险

不适合在 Callback 内做耗时很长的工作，因为该 Stream 后续操作要等它返回

另外，Stream 可以带有优先级：

int minPriority, maxPriority;

// Query the priority range for the device
cudaDeviceGetStreamPriorityRange(&minPriority, &maxPriority);

// Create two streams with different priorities
// cudaStreamDefault indicates the stream should be created with default flags
// in other words they will be blocking streams with respect to the legacy default stream
cudaStream_t stream1, stream2;
cudaStreamCreateWithPriority(&stream1, cudaStreamDefault, minPriority);  // Lowest priority
cudaStreamCreateWithPriority(&stream2, cudaStreamDefault, maxPriority);  // Highest priority

CUDA 中通常数值越小表示优先级越高，但应使用查询返回的范围，不要自行假设具体数字

优先级只是调度提示，不保证严格执行顺序，对内存传输不一定生效

Blocking, Non-blocking Streams and the Default Stream

如果启动 Kernel 时不指定 Stream，那么它会提交到 Default Stream，CUDA 提供了一个叫 Legacy Default Stream 的流，它在各个 Host Thread 之间共享

Legacy Default Stream 会和所有 Blocking Stream 相互同步，这是一个双向的隐式阻塞：

Legacy Default Stream 提交任务时，会等待所有已经发起的（Blocking）Stream 中的工作完成
Legacy Default Stream 提交任务时，所有随后的（Blocking）Stream 工作也必须等待它执行完毕才能继续

例如：

cudaStream_t stream1;
cudaStream_t stream2;

cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);

kernel1<<<grid, block, 0, stream1>>>(); // Run within Stream 1
kernel2<<<grid, block>>>();             // Run within Lagacy Default Stream
kernel3<<<grid, block, 0, stream2>>>(); // Run within Stream 2

即使三个 Kernel 没有数据依赖、硬件也有足够资源，它们仍可能被隐式串行化，kernel2 会等待 kernel1，kernel3 又会等待 kernel2，同样地，如果中间的 kernel2 换成一个 cudaMemcpy 也是这样（应该就是为这种场景设计的，让数据传输任务可以看见前面的任务，又对后面的任务可见）

只有一个例外，就是 Non-blocking Stream，它不与 Legacy Default Stream 发生隐式同步，这样创建：

cudaStream_t stream;
cudaStreamCreateWithFlags(&stream, cudaStreamNonBlocking);

// Create a non-blocking stream with priority
cudaStreamCreateWithPriority(&stream, cudaStreamNonBlocking, priority);

如果变成这样：

cudaStream_t stream1;
cudaStream_t stream2;

cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);

cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);

kernel1<<<grid, block, 0, stream1>>>();
kernel2<<<grid, block>>>();
kernel3<<<grid, block, 0, stream2>>>();

cudaDeviceSynchronize();

那么三个 Kernel 就可以并发执行了

注意，Blocking Stream 和 Non-blocking Stream 并不影响是否阻塞 Host

从 CUDA 7 开始，CUDA 允许每个 Host Thread 拥有自己的 Default Stream 而不是共享 Legacy Default Stream，这叫做 Per-thread Default Stream

这个特性必须显式开启（例如 NVCC 编译选项），开启后不会像 Legacy Default Stream 那样进行隐式同步，所以如果开了这个特性，Blocking Stream 和 Non-blocking Stream 表现就完全一样了

CUDA Event

Event 可以理解为插入 Stream 中的一个标记

它可以记录时间（使用 cudaEventElapsedTime(...) 函数）：

cudaStream_t stream;
cudaStreamCreate(&stream);

cudaEvent_t start;
cudaEvent_t stop;

// create the events
cudaEventCreate(&start);
cudaEventCreate(&stop);

// record the start event
cudaEventRecord(start, stream);

// launch the kernel
kernel<<<grid, block, 0, stream>>>(...);

// record the stop event
cudaEventRecord(stop, stream);

// wait for the stream to complete
// both events will have been triggered
cudaStreamSynchronize(stream);

// get the timing
float elapsedTime;
cudaEventElapsedTime(&elapsedTime, start, stop);
std::cout << "Kernel execution time: " << elapsedTime << " ms" << std::endl;

// clean up
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaStreamDestroy(stream);

这个功能可以用 cudaEventCreateWithFlags(&event, cudaEventDisableTiming) 的方式关闭

像 Stream 一样，它也可以阻塞等待（使用 cudaEventSynchronize(event) 函数）或者轮询（使用 cudaEventQuery(event) 函数）

Event 最重要的用途之一就是可以建立跨 Stream 依赖，比如 Stream B 依赖 Stream A 某个步骤生成的数据，可以这样干：

Stream A：… -> Produce data -> Record event -> …
Stream B：… -> Wait event -> Consume data -> …

即：

previousOperationsOfStreamA<<<..., streamA>>>();
producer<<<..., streamA>>>(buffer);
cudaEventRecord(ready, streamA);
otherOperationsOfStreamA<<<..., streamA>>>();

previousOperationsOfStreamB<<<..., streamB>>>();
cudaStreamWaitEvent(streamB, ready, 0);
consumer<<<..., streamB>>>(buffer);
otherOperationsOfStreamB<<<..., streamB>>>();

注意，这里的 cudaStreamWaitEvent(event) 并不会阻塞 CPU

Asynchronous Error Handling

常见检查方式是：

cudaPeekAtLastError() 读取但不清除最后错误
cudaGetLastError() 会读取并清除

// Some work occurs in streams.
cudaStreamSynchronize(stream);

// Look at the last error but do not clear it
cudaError_t err = cudaPeekAtLastError();
if (err != cudaSuccess) {
  printf("Error with name: %s\n", cudaGetErrorName(err));
  printf("Error description: %s\n", cudaGetErrorString(err));
}

// Look at the last error and clear it
cudaError_t err2 = cudaGetLastError();
if (err2 != cudaSuccess) {
  printf("Error with name: %s\n", cudaGetErrorName(err2));
  printf("Error description: %s\n", cudaGetErrorString(err2));
}

if (err2 != err) {
  printf("As expected, cudaPeekAtLastError() did not clear the error\n");
}

// Check again
cudaError_t err3 = cudaGetLastError();
if (err3 == cudaSuccess) {
  printf("As expected, cudaGetLastError() cleared the error\n");
}

调试异步错误时，可以临时设置：CUDA_LAUNCH_BLOCKING=1，它会让 Kernel Launch 后进行同步，更容易定位出错位置，但会显著降低运行速度

CUDA Graph

流式顺序执行的，但是有时需要一个 DAG，CUDA Graph 提供了这个功能

CUDA Graph 可以：

捕获一次操作序列或依赖 DAG，在图首次执行时构建或者手动构建
实例化出 graph 对象，这一步会建立执行图所需的所有运行时结构，从而尽可能加快图组件的启动速度
后续反复 Launch 已实例化的 graph，由于执行图操作所需的所有运行时结构均已就绪，因此图执行过程中的 CPU 开销降到了最低

例子：

constexpr int N = 500000 // tuned such that kernel takes a few microseconds

// A very lightweight kernel
__global__ void shortKernel(float* out_d, float* in_d) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if(idx < N) out_d[idx] = 1.23 * in_d[idx];
}

bool graphCreated = false;
cudaGraph_t graph;
cudaGraphExec_t instance;

// The graph will be executed NSTEP times
for(int istep = 0; istep < NSTEP; istep++) {
  if(!graphCreated) {
    // Capture the graph
    cudaStreamBeginCapture(stream, cudaStreamCaptureModeGlobal);

    // Launch NKERNEL kernels
    for(int ikrnl = 0; ikrnl < NKERNEL; ikrnl++){
      shortKernel<<<blocks, threads, 0, stream>>>(out_d, in_d);
    }

    // End the capture
    cudaStreamEndCapture(stream, &graph);

    // Instantiate the graph
    cudaGraphInstantiate(&instance, graph, NULL, NULL, 0);
    graphCreated=true;
  }

  // Launch the graph
  cudaGraphLaunch(instance, stream);

  // Synchronize the stream
  cudaStreamSynchronize(stream);
}

Unified and System Memory

Unified Virtual Address Space：CUDA 进程中的 CPU 内存和各 GPU 显存都位于同一个虚拟地址空间里，但每个处理器的内存仍对应不同的地址区间

可以使用 cudaPointerGetAttributes() 判断指针指向的是 CPU 内存还是某块 GPU 内存
cudaMemcpy*()（例如 cudaMemcpy()、cudaMemcpy2DAsync()）在传入 cudaMemcpyDefault 时可以自动判断复制方向

Unified Memory：允许同一块 Managed Memory 同时被 CPU 代码和 GPU Kernel 访问。所有 CUDA 支持的系统都支持 Unified Memory，但是分为几个层级：

Limited Unified Memory：

Windows、WSL 上（cudaDevAttrConcurrentManagedAccess=0）
只有显式地作为 Managed Memory 申请的内存才会被统一（也就是说只有 cudaMallocManaged、__managed__ 类似的方法申请的内存才会被视为统一内存）
Managed Memory 首先分配到 CPU 内存中，当 GPU 开始执行时，Managed Memory 会迁移到 GPU，GPU 活动期间，CPU 不得访问托管内存，当 GPU 同步时，托管内存会迁移回 CPU
不支持 GPU 侧按需细粒度迁移，Kernel 启动时，通常需要把所有 Managed Memory 迁移到 GPU，以避免 GPU 访问时发生无法处理的缺页
不允许超额分配（比如 GPU 显存只有 8G，那么不能申请 12G 的 Managed Memory）

Full Support for Explicit Managed Memory Allocations：

大多数 Linux 支持（cudaDevAttrConcurrentManagedAccess=1、cudaDevAttrPageableMemoryAccess=0）
也是只能显式使用类似 cudaMallocManaged 的方法进行分配
Managed Memory 通常分配在首次访问它的 CPU 或 GPU 的内存空间中，被其他处理器使用时通常会进行迁移
迁移或访问以内存页（Software Coherence）或 Cache Line（Hardware Coherence）为粒度
可以进行超额分配

Full Support for All Allocations With Software Coherence：

Linux + PCIe（cudaDevAttrConcurrentManagedAccess=1、cudaDevAttrPageableMemoryAccess=1、cudaDevAttrPageableMemoryAccessUsesHostPageTables=0）
以下内存都可以按 Unified Memory 的方式工作：malloc(...)、new ...、mmap(...) 和 cudaMallocManaged(...)
迁移或访问以内存页为粒度，可以超额分配
CPU 和 GPU 有各自的页表，由 Heterogeneous Memory Management（HMM）系统维护，这是由 Linux 内核提供的一套异构内存管理机制，主要通过缺页异常和页面迁移维护一致性，Linux 内核和 NVIDIA 驱动需要追踪当前页表在哪个设备的内存中，在必要时进行页表的迁移、失效

Full Support for All Allocations With Hardware Coherence：

NVLink C2C，如 Grace Hopper、Grace Blackwell（cudaDevAttrConcurrentManagedAccess=1、cudaDevAttrPageableMemoryAccess=1、cudaDevAttrPageableMemoryAccessUsesHostPageTables=1）
仍然是任何内存分配都被看做统一内存，可以超额分配
迁移或访问以 Cache Line 为粒度
通过 Address Translation Services（ATS）进行地址翻译即可，不需要 CPU 和 GPU 各自维护自己的页表，一致性冲突主要在发生共享的 Cache Line 上处理
支持 CPU–GPU 原生原子操作

还有一种叫做 Mapped Memory 的东西

在有 HMM 或 ATS 的系统中，所有 Host Memory 都可以通过原始的 Host Pointer 由 GPU 访问

然而在没有这两个的系统中，普通 malloc() 内存默认不能直接传给 Kernel，必须先将它注册并映射到 GPU 地址空间，此时：

cudaHostRegister(host_ptr, size, 0);

cudaHostRegister() 会把一段已经存在的 Host Memory 注册并页锁定

然后获取 GPU 端可以使用的地址，此时会转换成 GPU 虚拟空间下的 Device Pointer：

cudaHostGetDevicePointer(&device_ptr, host_ptr, 0);

Kernel 使用 device_ptr：

kernel<<<...>>>(device_ptr);

此时，数据没有先复制到显存，而是 GPU 每次读写时通过 PCIe 或 NVLink 跨互连访问

对 Mapped Host Memory 执行的 GPU 原子操作，从 CPU 或其他 GPU 的视角看，不保证是原子的

它适合少量、低频、一次性或流式访问，不适合作为大多数 Kernel 的主要内存

The NVIDIA CUDA Compiler（NVCC）

Parallel Thread Execution（PTX）：一种面向 GPU 的高级汇编语言，针对某种虚拟 GPU 架构，不同的设备可能只兼容特定版本的 PTX
Cubin：针对特定真实 GPU 架构的 SM 版本（例如 sm_120）的二进制格式
Fabin：可执行文件（会包含 CPU 代码和 GPU 代码）中的 GPU 代码保存的容器，为了兼容不同的 SM 版本，通常会包含几组 Cubin 二进制和一组或多组 PTX

The NVIDIA CUDA Compiler（NVCC）是用来编译 PTX、CUDA C/C++ 的工具链，属于 CUDA Toolkit 的一部分，包含编译器、链接器以及 PTX 和 Cubin 汇编器等多种工具

File Extension	Description	Content
`.c`	C source file	Host-only code
`.cpp`, `.cc`, `.cxx`	C++ source file	Host-only code
`.h`, `.hpp`, `.hh`, `.hxx`	C/C++ header file	Device code, host code, mix of host/device code
`.cu`	CUDA source file	Device code, host code, mix of host/device code
`.cuh`	CUDA header file	Device code, host code, mix of host/device code

NVCC 会区分出来 Host Code 和 Device Code，对于 Host Code 一般来说会调用普通的 C/C++ 编译器（gcc、clang），而对于 Device Code 会调用自己的 GPU 编译器

与普通的编译器相似，GPU Compiler 也会经历“代码 -> 汇编 -> 二进制”的阶段，然后再进行链接

例如对于下面的代码：

// ----- example.cu -----

#include <stdio.h>

__global__ void kernel() {
  printf("Hello from kernel\n");
}

void kernel_launcher() {
  kernel<<<1, 1>>>();
  cudaDeviceSynchronize();
}

int main() {
  kernel_launcher();
  return 0;
}

编译流程大概是这样：

High-level nvcc flow multiple architectures

NVCC 的 GPU 编译器也支持分离编译，但是它默认并不开启，即遵循 whole-program compilation，假定当前编译单元中能够找到其需要的全部设备代码，需要手动开启。同样的，为了弥补分离编译带来的跨文件优化损失，也可以使用链接时优化 LTO

NVCC 默认会压缩可执行文件或库中的 Fatbin，可以通过选项调整压缩策略（speed、size、balance、none）