CUDA 学习笔记 01 —— CUDA 基础

0. 学习 CUDA 的目的

作为一个算法工程师，平时接触 HPC (High Performance Computing) 的机会并不多，那为什么还要学习 CUDA 呢？
学习 CUDA 的目的不是为了用 CUDA 做模型加速，而是从 CUDA 角度理解目前较新的大模型设计理念，这些高性能模型是如何从原理上做到又快又好的。
例如火出圈的 DeepSeek 系列模型，在模型设计角度做了较多创新，并开源了部分 CUDA 代码，对于不了解 CUDA 的工程师，很难 get 到算法设计的优雅之处。
反观某家大模型基座公司，曾开源某个模型结构，论文中一通自夸，分析理论计算量有多低。但很多人实测发现速度并没有很快，究其原因，实际上是这家公司还用的小模型时代的旧思维，即：一个模型理论计算量低，那就是快。
大模型时代不了解硬件，不尊重硬件，在算法创新上不太可能走的远。

1. Hello World

cuda 代码

#include <cuda_runtime_api.h>
#include <stdio.h>
// cuda 中 host 表示 cpu 端，device 表示 gpu 端
// __device__ 是设备函数的声明符号，表明该函数在 device 执行，且只能在 device
// 中调用
__device__ const char *device_hello_world(void) {
  return "GPU: Hello world!\n";
}
// __host__ 是主机函数的声明符号，表明该函数在 host 执行，且只能在 host 中调用
__host__ const char *host_hello_world(void) { return "CPU: Hello world!\n"; }
// __global__ 是核函数的声明符号，表明该函数在 device 执行，且只能在 host 中调用
__global__ void hello_world(void) {
  const char *str = device_hello_world();
  printf("%s", str);
}
int main(int argc, char **argv) {
  printf("%s", host_hello_world());
  // <<<grid_dim, block_dim>>> 是核函数的调用符号，表示启动 grid_dim 个 block，
  // 每个 block 有 block_dim 个线程
  hello_world<<<1, 10>>>();
  cudaDeviceReset();
  return 0;
}

cuda 的三个函数声明符号：
- __host__：主机函数，表示该函数在 CPU 上执行，且只能在 CPU 中调用
- __device__：设备函数，表示该函数在 GPU 上执行，且只能在 GPU 中调用
- __global__：核函数，表示该函数在 GPU 上执行，且只能在 CPU 中调用
其中 __global__ 声明的函数类型被称为 核函数，是 CUDA 中最重要的函数类型
- 核函数通过 <<<grid_dim, block_dim>>> 的方式调用，其中 <<<>>> 是 cuda 扩展关键字
- grid_dim 表示启动的 block 数量，block_dim 表示每个 block 中的线程数量
- grid_dim 和 block_dim 都是 dim3 类型的变量，表示三维数组，如果使用整形则模型 y 和 z 维度都为 1

编译

1	nvcc hello_world.cu -g -o hello_world

运行

./hello_world
# CPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!
# GPU: Hello world!

2. dimension 测试

#include <cuda_runtime.h>
#include <stdio.h>
__global__ void checkIndex(void) {
  printf("threadIdx:(%d,%d,%d) blockIdx:(%d,%d,%d) blockDim:(%d,%d,%d)\
  gridDim(%d,%d,%d)\n",
         threadIdx.x, threadIdx.y, threadIdx.z, blockIdx.x, blockIdx.y,
         blockIdx.z, blockDim.x, blockDim.y, blockDim.z, gridDim.x, gridDim.y,
         gridDim.z);
}
int main(int argc, char **argv) {
  int nElem = 6;  // number of elements
  dim3 block(3);  // block size
  int nBlock = (nElem + block.x - 1) / block.x; // number of blocks
  dim3 grid(nBlock);  // grid size
  printf("grid.x %d grid.y %d grid.z %d\n", grid.x, grid.y, grid.z);  
  printf("block.x %d block.y %d block.z %d\n", block.x, block.y, block.z);
  checkIndex<<<grid, block>>>();
  cudaDeviceReset();
  return 0;
}

执行结果：

./check_dimension
# grid.x 2 grid.y 1 grid.z 1
# block.x 3 block.y 1 block.z 1
# threadIdx:(0,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(1,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(2,0,0) blockIdx:(1,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(0,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(1,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)
# threadIdx:(2,0,0) blockIdx:(0,0,0) blockDim:(3,1,1)  gridDim(2,1,1)

3. CUDA 向量加法

代码

#include "freshman.h"
#include <cuda_runtime.h>
#include <stdio.h>
__host__ void sumArrays(float *a, float *b, float *res, const int size) {
  for (int i = 0; i < size; i++) {
    res[i] = a[i] + b[i];
  }
}
__global__ void sumArraysGPU(float *a, float *b, float *res, const int size) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < size) // 线程索引越界检查
    res[i] = a[i] + b[i];
}
int main(int argc, char **argv) {
  // set up device
  initDevice(0);
  // allocate host memory
  int nElem = 1 << 24;
  printf("Vector size:%d\n", nElem);
  int nByte = sizeof(float) * nElem;
  float *a_h = (float *)malloc(nByte);
  float *b_h = (float *)malloc(nByte);
  float *res_h = (float *)malloc(nByte);
  float *res_from_gpu_h = (float *)malloc(nByte);
  memset(res_h, 0, nByte);
  memset(res_from_gpu_h, 0, nByte);
  // allocate device memory
  float *a_d, *b_d, *res_d;
  CHECK(cudaMalloc((float **)&a_d, nByte));
  CHECK(cudaMalloc((float **)&b_d, nByte));
  CHECK(cudaMalloc((float **)&res_d, nByte));
  // randomly initialize the input data
  initialData(a_h, nElem);
  initialData(b_h, nElem);
  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));
  // set up execution configuration
  // 1. 计算最佳块大小
  int minGridSize, bestBlockSize;
  cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize,
                                     (void *)sumArraysGPU,
                                     0, // 动态共享内存大小
                                     0  // 无块大小限制
  );
  // 2. 设置网格和块维度
  dim3 block(bestBlockSize);
  dim3 grid((nElem + bestBlockSize - 1) / bestBlockSize);
  // 3. 设备执行并统计耗时
  double iStart, iElaps;
  iStart = cpuSecond();
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d, nElem);
  iElaps = cpuSecond() - iStart;
  CHECK(cudaGetLastError());
  CHECK(cudaDeviceSynchronize());
  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
  printf("Execution configuration<<<%d,%d>>> Time elapsed %f sec\n", grid.x,
         block.x, iElaps);
  // 4. CPU执行并统计耗时
  iStart = cpuSecond();
  sumArrays(a_h, b_h, res_h, nElem);
  iElaps = cpuSecond() - iStart;
  printf("CPU Time elapsed %f sec\n", iElaps);
  // 5. 检查结果
  checkResult(res_h, res_from_gpu_h, nElem);
  // 6. 释放内存
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);
  free(a_h);
  free(b_h);
  free(res_h);
  free(res_from_gpu_h);
  return 0;
}

cudaOccupancyMaxPotentialBlockSize 函数用于计算最佳块大小
max thread per block 是 cuda 中的一个限制，表示每个块中最多可以有多少个线程，一般为 1024，当超过这个限制时，CHECK(cudaGetLastError()); 会报错

运行结果

./sum_arrays
# Using device 0: NVIDIA GeForce RTX 4060 Laptop GPU
# Vector size:16777216
# Execution configuration<<<21846,768>>> Time elapsed 0.000030 sec
# CPU Time elapsed 0.076604 sec
# Check result success!

可以看出，cuda 的执行时间远远小于 cpu 的执行时间，相差了 2553 倍

总体流程

graph TD
    A[Host 端申请内存] --> B[Host 端初始化输入数据]
    B --> C[Device 端申请内存]
    C --> D[拷贝 Host 输入数据到 Device 端]
    D --> E[Device 端执行核函数]
    E --> F[拷贝 Device 端输出数据到 Host 端]
    F --> G[Host 端检查结果]
    G --> H[释放 Host 端和 Device 端全部内存]
    B --> J[Host 端执行普通函数] --> G

细节分析

1. `cuda` 内存分配是怎么做的？

1 2	float a_d; // 空指针 cudaMalloc((float *)&a_d, nByte); // 将指针的地址转成二级指针（指针的指针）传入内存分配函数

这里的二级指针应用很巧妙，由于 c++ 中的指针是值传递，所以如果是一级指针传入 cudaMalloc 函数时，指针 a_d 的值不会改变，因此只能将指针的地址转成二级指针传入内存分配函数

2. `cuda` 内存拷贝是怎么做的？

1	cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice);

cudaMemcpy 函数的四个参数分别是：
- dst：目标地址
- src：源地址
- size：拷贝的字节数
- kind：拷贝的类型，属于 cudaMemcpyKind 枚举类型
  - cudaMemcpyHostToHost = 0, /**< Host -> Host */
  - cudaMemcpyHostToDevice = 1, /**< Host -> Device */
  - cudaMemcpyDeviceToHost = 2, /**< Device -> Host */
  - cudaMemcpyDeviceToDevice = 3, /**< Device -> Device */
  - cudaMemcpyDefault = 4 /**< Direction of the transfer is inferred from the pointer values. Requires unified virtual addressing */

3. `cuda` 核函数如何解析线程索引？

1
2
3

int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < size) // 线程索引越界检查
  res[i] = a[i] + b[i];

blockIdx：表示当前块的索引
blockDim：表示当前块的维度（每个块中的线程数）
threadIdx：表示当前线程的索引
每个线程中计算两个标量的和
由于 gridDim * blockDim 可能大于 size，所以需要判断线程索引是否越界

4. 如何计算最佳块大小？

int minGridSize, bestBlockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &bestBlockSize,
                                     (void *)sumArraysGPU,
                                     0, // 动态共享内存大小
                                     0  // 无块大小限制
);

cudaOccupancyMaxPotentialBlockSize 函数用于计算最佳块大小，五个参数分别是：
- minGridSize：最小网格大小变量地址
- bestBlockSize：最佳块大小变量地址
- kernel：核函数指针
- dynamicSMemSize：动态共享内存大小
- blockSizeLimit：块大小限制
函数名就是函数地址，可强转为 void * 函数指针（也可以写成：(void *)&sumArraysGPU）

4. CUDA 编程模型

线程块

线程块 block 是 CUDA 中的逻辑执行单元，是一个三维逻辑结构：
- block.x：表示块的 x 维度大小
- block.y：表示块的 y 维度大小
- block.z：表示块的 z 维度大小
- 其中 block.x 是最内层的循环，block.y 是第二层循环，block.z 是最外层的循环
- 用三维数组可以表示为：tread[z][y][x]，即 tid = threadIdx.x + threadIdx.y * blockDim.x + threadIdx.z * blockDim.x * blockDim.y

线程束

线程束 warp 是 CUDA 基本调度执行单元，一个 warp 由 32 个线程组成
- 一个 warp 中的线程在一个时钟周期内执行同一条指令（单指令多线程，SIMT）
- 一个 warp 中的线程可以共享指令指针和执行资源（如寄存器、缓存等）
- Warp 调度器（warp scheduler）负责将 warp 分配到物理执行单元上执行
线程块会被划分为多个 warp， $WarpsPerBlock=ceil(\frac{ThreadsPerBlock}{32})$

CUDA core

CUDA core 是 CUDA 物理执行单元，负责实际的计算任务
一个 CUDA core 一个时钟周期只能计算一个线程的指令

StreamMultiprocessor

StreamMultiprocessor 流式多处理器（简称 SM），负责执行 CUDA 线程块中的并行计算任务
每个 GPU 包含多个 SM，每个 SM 包含多个 CUDA core，例如：RTX 4060 有 24 个 SM，每个 SM 有 128 个 CUDA core

5. Reduce

规约（Reduce）是 CUDA 编程中常见的操作，主要用于将多个数据元素规约为一个数据元素
规约操作通常是一个二元操作，例如：sum、mul、max、min 等，简单的规约可以合并成强大的算子，甚至可以说规约算子是神经网络的基础

规约求和

// CPU 规约求和
int recursiveReduce(int *data, int const size) {
  // terminate check
  if (size == 1)
    return data[0];
  // renew the stride
  int const stride = size / 2;
  if (size % 2 == 1) {
    for (int i = 0; i < stride; i++) {
      data[i] += data[i + stride];
    }
    data[0] += data[size - 1];
  } else {
    for (int i = 0; i < stride; i++) {
      data[i] += data[i + stride];
    }
  }
  // call
  return recursiveReduce(data, stride);
}
// GPU 规约相邻求和
__global__ void reduceNeighbored(int *g_idata, int *g_odata, unsigned int n) {
  // set thread ID
  unsigned int tid = threadIdx.x;
  // boundary check
  if (tid >= n)
    return;
  // convert global data pointer to local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  // in-place reduction in global memory
  for (int stride = 1; stride < blockDim.x; stride *= 2) {
    if ((tid % (2 * stride)) == 0) {
      idata[tid] += idata[tid + stride];
    }
    // synchronize within block
    __syncthreads();
  }
  // write result for this block to global mem
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}
// GPU 规约相邻求和（简化版）
__global__ void reduceNeighboredLess(int *g_idata, int *g_odata,
                                     unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
  // convert global data pointer to the local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  if (idx > n)
    return;
  // in-place reduction in global memory
  for (int stride = 1; stride < blockDim.x; stride *= 2) {
    // convert tid into local array index
    int index = 2 * stride * tid;
    if (index < blockDim.x) {
      idata[index] += idata[index + stride];
    }
    __syncthreads();
  }
  // write result for this block to global men
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}
// GPU 规约交错求和，主要是 stride 的计算方式不同
__global__ void reduceInterleaved(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
  // convert global data pointer to the local point of this block
  int *idata = g_idata + blockIdx.x * blockDim.x;
  if (idx >= n)
    return;
  // in-place reduction in global memory
  for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
      idata[tid] += idata[tid + stride];
    }
    __syncthreads();
  }
  // write result for this block to global men
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

效率分析

本代码中使用了三种核函数实现方式做同一个规约操作，分别是：
- reduceNeighbored：相邻线程规约
- reduceNeighboredLess：相邻线程规约（简化版）
- reduceInterleaved：交错线程规约
三者效率从高到低依次是：
- reduceInterleaved > reduceNeighboredLess > reduceNeighbored
三者的示意图分别如下：
- reduceNeighbored：相邻线程规约的实现
- reduceNeighboredLess：相邻线程规约的简化版实现（注意圆圈中的符号已和上图不一致）
- reduceInterleaved：交错线程规约的实现
三者效率差异主要来自于 线程分支分化，后续会详细介绍

6. 循环展开

循环展开（Loop Unrolling）是 CUDA 中常用的优化手段，主要用于减少循环控制开销和提高指令级并行度
简单说就是一个线程不再只计算一个数据，而是计算多个数据，而且是直接在代码中展开，而不是在编译器中展开
可以简单理解成：启动线程是需要花时间的，启动一个线程只算一个数据，太浪费了，所以我们可以让一个线程计算多个数据，这样就能减少启动线程的时间开销，所以就省时间了

代码

// 数据总长是 8 * blockDim.x * gridDim.x
// 线程数是 blockDim.x * gridDim.x
// 每个线程计算 8 个数据
__global__ void reduceUnroll8(int *g_idata, int *g_odata, unsigned int n) {
  unsigned int tid = threadIdx.x;
  unsigned int idx = blockDim.x * blockIdx.x * 8 + threadIdx.x;
  if (tid >= n)
    return;
  int *idata = g_idata + blockIdx.x * blockDim.x * 8;
  // 循环展开，每个线程计算 8 个数据
  // 直接把 8 * blockDim.x * gridDim.x 的数据总长
  // 聚合到了 blockDim.x * gridDim.x 的线程数上
  if (idx + 7 * blockDim.x < n) {
    g_idata[idx] += g_idata[idx + blockDim.x];
    g_idata[idx] += g_idata[idx + blockDim.x * 2];
    g_idata[idx] += g_idata[idx + blockDim.x * 3];
    g_idata[idx] += g_idata[idx + blockDim.x * 4];
    g_idata[idx] += g_idata[idx + blockDim.x * 5];
    g_idata[idx] += g_idata[idx + blockDim.x * 6];
    g_idata[idx] += g_idata[idx + blockDim.x * 7];
  }
  // 这里需要同步，也就是线程阻塞直到所有线程都执行完
  __syncthreads();
  // 然后就是一个最简单的规约操作了，和上面一样
  for (int stride = blockDim.x / 2; stride > 0; stride >>= 1) {
    if (tid < stride) {
      idata[tid] += idata[tid + stride];
    }
    __syncthreads();
  }
  // write result for this block to global mem
  if (tid == 0)
    g_odata[blockIdx.x] = idata[0];
}

7. 核函数递归调用

和 CPU 一样，CUDA 也支持核函数递归调用，调用方式和普通递归函数一样
需要注意的是在编译的时候需要加上 -rdc=true 选项

代码

__global__ void nesthelloworld(int iSize, int iDepth) {
  unsigned int tid = threadIdx.x;
  printf("depth : %d blockIdx: %d,threadIdx: %d\n", iDepth, blockIdx.x,
         threadIdx.x);
  if (iSize == 1)
    return;
  int nthread = (iSize >> 1);
  if (tid == 0 && nthread > 0) {
    // 递归调用核函数
    nesthelloworld<<<1, nthread>>>(nthread, ++iDepth);
    printf("-----------> nested execution depth: %d\n", iDepth);
  }
}

8. 固定内存

Pinned Memory 是 CUDA 中的一种特殊内存类型（不是显存，是内存），主要用于提高数据传输效率
普通内存是分页管理，分页管理存在两个问题：
1. 一页内存逻辑上连续，但物理上不连续
2. 操作系统可能会将内存页交换到磁盘上，导致数据不在物理内存中
Pinned Memory 就是解决了这两个问题，分配了一块连续物理地址且固定的主机内存（host 内存），方便整块拷贝数据到显存（DMA）

代码

#include "freshman.h"
#include <cuda_runtime.h>
int main(int argc, char **argv) {
  int dev = 0;
  cudaSetDevice(dev);
  int nElem = 1 << 14;
  int nByte = sizeof(float) * nElem;
  float *a_h, *b_h, *res_h, *res_from_gpu_h;
  // 注意这里的 cudaMallocHost 和 cudaMalloc 是不同的
  // 前者申请的是 host 固定内存，后者申请的是 device 显存
  // cudaMallocHost 是 malloc 的一个平替
  CHECK(cudaMallocHost((float **)&a_h, nByte));
  CHECK(cudaMallocHost((float **)&b_h, nByte));
  CHECK(cudaMallocHost((float **)&res_h, nByte));
  CHECK(cudaMallocHost((float **)&res_from_gpu_h, nByte));
  // 初始化数据
  memset(res_h, 0, nByte);
  memset(res_from_gpu_h, 0, nByte);
  initialData(a_h, nElem);
  initialData(b_h, nElem);
  // 申请设备显存
  float *a_d, *b_d, *res_d;
  CHECK(cudaMalloc((float **)&a_d, nByte));
  CHECK(cudaMalloc((float **)&b_d, nByte));
  CHECK(cudaMalloc((float **)&res_d, nByte));
  // 拷贝数据到设备显存
  CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
  CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));
  // 跑核函数
  dim3 block(1024);
  dim3 grid(nElem / block.x);
  sumArraysGPU<<<grid, block>>>(a_d, b_d, res_d);
  printf("Execution configuration<<<%d,%d>>>\n", grid.x, block.x);
  // 结果拷贝回主机并检查两个设备计算结果是否一致
  CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
  sumArrays(a_h, b_h, res_h, nElem);
  checkResult(res_h, res_from_gpu_h, nElem);
  // 释放内存
  // 注意这里的 cudaFreeHost 和 cudaFree 是不同的
  // 前者释放的是 host 固定内存，后者释放的是 device 显存
  cudaFree(a_d);
  cudaFree(b_d);
  cudaFree(res_d);
  cudaFreeHost(a_h);
  cudaFreeHost(b_h);
  cudaFreeHost(res_h);
  cudaFreeHost(res_from_gpu_h);
  return 0;
}

关键点

主存 普通内存 的分配和释放函数是 malloc 和 free
主存 固定内存 的分配和释放函数是 cudaMallocHost 和 cudaFreeHost
显存的分配和释放函数是 cudaMalloc 和 cudaFree

9. 零拷贝内存和统一虚拟地址

Zero-Copy Memory 是 CUDA 中一种允许 GPU 直接访问主机内存的技术，避免了显式的数据拷贝操作（不需要用 cudaMemcpy 函数）
实际上，Zero-Copy Memory 在很多时候并不快，因为 GPU 访问主机内存的速度远远低于访问显存的速度，因此，Zero-Copy Memory 只适用于一些特殊的场景，例如：
- 主机内存中的数据只需要被 GPU 使用一次
- 数据量太大，显存放不下
- 调试用途
Zero-Copy Memory 的实现方式是将主机内存映射到 GPU 的地址空间中，GPU 通过访问这个地址空间来访问主机内存，实际上走的是 PCIe 总线
由于不需要先完成所有数据的拷贝再开始执行核函数，因此 Zero-Copy Memory 使用异步拷贝的方式来实现，可将部分拷贝数据的时间和核函数执行的时间重叠，但并不多
Unified Virtual Addressing (UVA) 是 CUDA 中的一种内存管理机制，允许 CPU 和 GPU 共享同一虚拟地址空间

代码

float *a_host, *b_host, *res_d;
// 申请主机固定内存，添加特殊 flag cudaHostAllocMapped
CHECK(cudaMallocHost((float **)&a_host, nByte, cudaHostAllocMapped));
CHECK(cudaMallocHost((float **)&b_host, nByte, cudaHostAllocMapped));
// a_host 和 b_host 是可直接作为核函数的输入参数
// 也可以转成 Device 地址空间，如下：
float *a_dev, *b_dev;
// 映射主机内存到设备地址空间
CHECK(cudaHostGetDevicePointer((void **)&a_dev, (void *)a_host, 0));
CHECK(cudaHostGetDevicePointer((void **)&b_dev, (void *)b_host, 0));
// 用 a_dev 和 b_dev 作为核函数的输入参数计算

和 Zero-Copy Memory 一样，UVA 甚至不需要将 a_host 转成 a_dev，直接用 a_host 就可以调用核函数

10. Aos 和 SoA

CUDA 不仅支持最简单的原生数据类型，还支持自定义数据类型（struct），例如 Aos 和 SoA 等
Aos（Array of Structures）和 SoA（Structure of Arrays）是两种不同的数据存储方式，这两种方式由于变量的排布方式不同，导致了访问内存的效率差异

Aos

Aos 是将多个结构体存储在一个数组中，每个结构体的成员变量是连续存储的
例如：

struct AoSStruct {
  float a;
  float b;
};
__global__ void sumArraysGPU(float *a, float *b, struct naiveStruct *res, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    res[i].a = a[i] + b[i];
}

这里的 res 是一个 结构体数组，a 和 b 是结构体的成员变量，每个变量都是一个 float 类型的标量
每个结构体的成员变量是连续存储的，即：a1 b1 a2 b2 a3 b3 ...

SoA

SoA 是将多个结构体的成员变量存储在一个数组中，每个成员变量是连续存储的
例如：

struct SoAStruct {
  float a[SIZE];
  float b[SIZE];
};
__global__ void sumArraysGPU(float *a, float *b, struct SoAStruct *res, int n) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  if (i < n)
    res->a[i] = a[i] + b[i];
}

这里的 res 是一个 结构体，a 和 b 是结构体的成员变量，每个变量都是一个 float 类型的数组
每个成员变量的数组是连续存储的，即：a1 a2 a3 ... b1 b2 b3 ...

11. 行主序和列主序

行主序（Row Major Order）和 列主序（Column Major Order）是两种不同的数组存储方式
行主序 是将数组的每一行存储在连续的内存中，列主序 是将数组的每一列存储在连续的内存中
例如：

int a[3][4] = {
    {1, 2, 3, 4},
    {5, 6, 7, 8},
    {9, 10, 11, 12}};

行主序 存储方式是：1 2 3 4 5 6 7 8 9 10 11 12
列主序 存储方式是：1 5 9 2 6 10 3 7 11 4 8 12
默认情况下，C / C++ / CUDA 语言是 行主序 存储方式
在行主序存储下，如果按行序访问数组元素，访问效率会更高，因为连续的内存访问会提高缓存命中率，反之如果按列序访问数组元素，访问效率会更低

速度对比

WIP

0. 学习 CUDA 的目的

1. Hello World

cuda 代码

编译

运行

2. dimension 测试

3. CUDA 向量加法

代码

运行结果

总体流程

细节分析

1. cuda 内存分配是怎么做的？

2. cuda 内存拷贝是怎么做的？

3. cuda 核函数如何解析线程索引？

4. 如何计算最佳块大小？

4. CUDA 编程模型

线程块

线程束

CUDA core

StreamMultiprocessor

5. Reduce

规约求和

效率分析

6. 循环展开

代码

7. 核函数递归调用

代码

8. 固定内存

代码

关键点

9. 零拷贝内存 和 统一虚拟地址

代码

10. Aos 和 SoA

Aos

SoA

11. 行主序和列主序

速度对比

1. `cuda` 内存分配是怎么做的？

2. `cuda` 内存拷贝是怎么做的？

3. `cuda` 核函数如何解析线程索引？

9. 零拷贝内存和统一虚拟地址