第2章 CUDA 编程基础 - Hello World on GPU（Getting Started with CUDA）

一、导读

1.1 本章学习目标

通过本章的学习，你将能够：

搭建完整的 CUDA 开发环境并验证安装正确性

理解 CUDA 程序的基本结构和执行流程

掌握主机代码与设备代码的区别及标识方法

编写并运行第一个 CUDA 内核函数

理解线程层次结构：线程、线程块、网格的概念

掌握 CUDA 内存管理的基本 API：cudaMalloc、cudaMemcpy、cudaFree

学会使用 nvcc 编译器编译 CUDA 程序

掌握基本的 CUDA 错误处理机制

深入理解 GPU 架构对编程模型的影响

能够诊断和解决常见的 CUDA 编程问题

1.2 为什么本章如此重要

本章是你 CUDA 编程之旅的起点。就像学习任何编程语言一样，第一个程序往往是"Hello World"。在 CUDA 中，这个简单的程序背后蕴含着丰富的概念：主机与设备的分离、内核函数的启动、线程的组织方式等。

理解这些基础概念至关重要，因为它们是后续所有高级特性的基石。如果你在这一章留下知识盲区，后续学习线程同步、内存优化等内容时会感到困惑。因此，请务必动手实践本章的每一个代码示例，确保真正理解每一行代码的含义。

CUDA 编程模型的核心思想是将计算任务分解为大量并行执行的线程。本章将向你展示如何定义这些线程、如何组织它们、如何让它们在 GPU 上协同工作。一旦你掌握了这一基本模式，你就具备了编写任何 CUDA 程序的能力。

此外，本章还将介绍 CUDA 开发工具链的基本使用方法。熟悉这些工具将大大提高你的开发效率和问题排查能力。

GPU 计算的革命性意义：

GPU 的出现改变了高性能计算的格局。传统 CPU 设计注重低延迟和通用性，而 GPU 则专注于高吞吐量并行计算。这种架构差异使得 GPU 在以下场景中具有显著优势：

大规模数据并行：图像处理、信号处理、科学模拟等需要对大量数据执行相同操作的场景

深度学习训练与推理：神经网络中的矩阵运算天然适合 GPU 并行执行

密码学与区块链：哈希计算、加密解密等操作可以从 GPU 加速中受益

金融建模：蒙特卡洛模拟、期权定价等需要大量随机采样的应用

理解 CUDA 编程基础是进入这些应用领域的第一步。

1.3 核心问题引导

在开始之前，请思考以下问题：

CUDA 程序与普通 C/C++ 程序在执行模型上有什么根本区别？

什么是"主机"（Host）和"设备"（Device）？它们之间如何通信？

内核函数（Kernel）是什么？如何启动一个内核函数？

线程块（Thread Block）和网格（Grid）的作用是什么？

为什么需要显式地在主机和设备之间传输数据？

nvcc 编译器如何处理 CUDA 代码中的主机部分和设备部分？

如何检测 CUDA API 调用是否成功？

第一个 CUDA 程序的执行流程是怎样的？

二、核心概念详解

2.1 主机与设备模型

CUDA 编程模型建立在异构计算的基础上，其中系统包含两种类型的处理器：

主机（Host）

主机指的是 CPU 及其关联的主内存。主机负责：

执行程序的控制逻辑
管理设备资源
准备和预处理数据
启动 GPU 内核函数
后处理 GPU 返回的结果

主机代码使用标准的 C/C++ 编译，运行在 CPU 上。

设备（Device）

设备指的是 GPU 及其显存。设备负责：

执行大规模并行计算
存储计算所需的数据和中间结果

设备代码使用 CUDA C/C++ 编写，经过 nvcc 编译后在 GPU 上运行。

主机-设备交互

主机和设备拥有各自独立的内存空间，这意味着：

主机不能直接访问设备内存
设备不能直接访问主机内存
数据必须通过显式的内存传输操作在两者之间移动

这种分离是 GPU 架构的自然结果：GPU 拥有独立的高速显存（GDDR6 或 HBM），与 CPU 主存物理隔离。数据传输通过 PCIe 总线进行，带宽远低于 GPU 内部内存带宽。

┌─────────────────────┐         ┌─────────────────────┐
│      Host (CPU)     │         │    Device (GPU)     │
│                     │         │                     │
│  ┌───────────────┐  │         │  ┌───────────────┐  │
│  │  Host Memory  │  │         │  │ Device Memory │  │
│  └───────┬───────┘  │         │  └───────┬───────┘  │
│          │          │         │          │          │
│  ┌───────▼───────┐  │         │  ┌───────▼───────┐  │
│  │  Host Code    │  │         │  │ Device Code   │  │
│  │  (C/C++)      │  │         │  │ (CUDA Kernel) │  │
│  └───────────────┘  │         │  └───────────────┘  │
└─────────────────────┘         └─────────────────────┘
              │                           │
              └────── PCIe Bus ───────────┘
                    cudaMemcpy()

理解这一模型对于编写高效的 CUDA 程序至关重要。频繁的数据传输会严重降低性能，因此优化的关键之一是尽量减少主机与设备之间的数据交换。

PCIe 带宽限制：

PCIe Gen3 x16 的理论带宽约为 16 GB/s，而现代 GPU 的显存带宽可达 500-900 GB/s。这意味着数据传输可能成为性能瓶颈。在实际应用中，应该：

尽量减少数据传输次数

使用异步传输重叠计算和通信

考虑使用统一内存（Unified Memory）简化数据管理

2.2 CUDA 程序的基本结构

一个典型的 CUDA 程序包含以下步骤：

分配主机内存：使用 malloc 或 new 分配 CPU 端的内存

初始化主机数据：填充输入数据

分配设备内存：使用 cudaMalloc 分配 GPU 端的内存

传输数据到设备：使用 cudaMemcpy 将输入数据从主机复制到设备

启动内核函数：使用 <<<grid, block>>> 语法启动 GPU 计算

传输结果回主机：使用 cudaMemcpy 将结果从设备复制回主机

释放设备内存：使用 cudaFree 释放 GPU 内存

处理和验证结果：在主机端检查结果的正确性

这种模式被称为"离载"（Offload）模式：计算密集型任务被卸载到 GPU 上执行，而控制逻辑保留在 CPU 上。

详细步骤解析：

步骤1-2：主机内存分配与初始化

cpp

int n = 1024;
float *h_data = (float *)malloc(n * sizeof(float));
for (int i = 0; i < n; i++) {
    h_data[i] = (float)i;
}

步骤3：设备内存分配

cpp

float *d_data;
cudaMalloc((void **)&d_data, n * sizeof(float));

注意：cudaMalloc 返回的是设备指针，只能在设备代码中使用。

步骤4：数据传输

cpp

cudaMemcpy(d_data, h_data, n * sizeof(float), cudaMemcpyHostToDevice);

步骤5：启动内核

cpp

myKernel<<<gridSize, blockSize>>>(d_data, n);

步骤6-7：获取结果与清理

cpp

cudaMemcpy(h_data, d_data, n * sizeof(float), cudaMemcpyDeviceToHost);
cudaFree(d_data);
free(h_data);

2.3 内核函数（Kernel）

内核函数是 CUDA 编程的核心抽象。它是一个在 GPU 上执行的函数，由大量并行线程同时执行。

定义内核函数

内核函数使用 __global__ 修饰符声明：

cpp

__global__ void myKernel(float *data, int n) {
    // 每个线程执行的代码
}

__global__ 告诉编译器这个函数：

可以从主机代码调用
在设备上执行
返回类型必须是 void

其他函数修饰符：

__device__：只能在设备上调用，在设备上执行
__host__：只能在主机上调用，在主机上执行（默认）
__global__：可以在主机上调用，在设备上执行

启动内核函数

内核函数使用特殊的执行配置语法启动：

cpp

myKernel<<<gridSize, blockSize>>>(arguments);

其中：

gridSize：网格中的线程块数量
blockSize：每个线程块中的线程数量

这两个参数可以是整数（一维）或 dim3 类型（多维）。

线程索引

在内核函数中，每个线程可以通过内置变量确定自己的身份：

threadIdx.x：线程在线程块内的索引
blockIdx.x：线程块在网格内的索引
blockDim.x：线程块的大小（线程数量）
gridDim.x：网格的大小（线程块数量）

全局线程索引的计算公式为：

cpp

int globalIdx = blockIdx.x * blockDim.x + threadIdx.x;

这个公式将二维的线程层次结构映射为一维的全局索引，用于访问数组元素。

多维线程配置：

对于处理图像、矩阵等多维数据，可以使用多维线程配置：

cpp

dim3 blockSize(16, 16);  // 256 线程每块
dim3 gridSize((width + 15) / 16, (height + 15) / 16);

__global__ void imageKernel(float *image, int width, int height) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (col < width && row < height) {
        int idx = row * width + col;
        image[idx] *= 2.0f;
    }
}

2.4 线程层次结构

CUDA 采用两级线程层次结构：

线程（Thread）

最基本的执行单元。每个线程执行相同的内核代码，但可以通过线程索引处理不同的数据。

线程块（Thread Block）

一组协作的线程。线程块内的线程可以：

通过共享内存进行通信
使用同步原语协调执行
共享相同的线程块索引

线程块的大小有限制，取决于 GPU 的计算能力。通常最大为 1024 个线程每块。

网格（Grid）

一组线程块。网格中的所有线程块执行相同的内核函数。网格的大小决定了总的并行度。

为什么需要两级结构？

两级结构的设计有以下考虑：

硬件映射：线程块映射到 GPU 的流多处理器（SM）上。每个 SM 可以同时执行多个线程块。

可扩展性：线程块是独立的执行单元，可以以任意顺序调度。这使得 CUDA 程序可以在不同规模的 GPU 上运行，无需修改代码。

同步范围：线程块内的线程可以同步，但不同线程块之间不能直接同步。这简化了硬件设计，同时满足了大多数并行算法的需求。

Warp 概念：

在每个线程块内，线程被进一步组织成 warp。一个 warp 包含 32 个线程，这些线程以 SIMD（单指令多数据）方式执行。也就是说，warp 中的所有线程在同一时钟周期执行相同的指令，但处理不同的数据。

理解 warp 对于性能优化至关重要：

如果 warp 中的线程出现分支发散（即不同线程走不同的分支），性能会下降
内存访问应该在 warp 级别合并，以最大化带宽利用率

2.5 内存管理 API

CUDA 提供了一套 C 风格的 API 用于设备内存管理：

cudaMalloc

cpp

cudaError_t cudaMalloc(void **devPtr, size_t size);

在设备上分配内存。返回的指针只能在设备代码中使用。

cudaMemcpy

cpp

cudaError_t cudaMemcpy(void *dst, const void *src, size_t count, 
                       enum cudaMemcpyKind kind);

在主机和设备之间复制数据。kind 参数指定传输方向：

cudaMemcpyHostToDevice：主机到设备
cudaMemcpyDeviceToHost：设备到主机
cudaMemcpyDeviceToDevice：设备到设备
cudaMemcpyHostToHost：主机到主机

cudaFree

cpp

cudaError_t cudaFree(void *devPtr);

释放设备上分配的内存。

cudaMemset

cpp

cudaError_t cudaMemset(void *devPtr, int value, size_t count);

将设备内存设置为指定值。

异步内存传输：

除了同步的 cudaMemcpy，CUDA 还提供异步版本：

cpp

cudaMemcpyAsync(void *dst, const void *src, size_t count, 
                enum cudaMemcpyKind kind, cudaStream_t stream);

异步传输立即返回，允许主机继续执行其他任务。这对于重叠计算和通信非常有用。

2.6 nvcc 编译器

nvcc 是 NVIDIA 提供的 CUDA 编译器驱动程序。它的工作流程如下：

预处理：处理宏和包含文件

分离：将代码分为主机部分和设备部分

主机编译：使用系统 C/C++ 编译器（如 gcc、cl）编译主机代码

设备编译：将设备代码编译为 PTX（Parallel Thread Execution）中间表示

链接：将主机目标代码和设备二进制代码链接为可执行文件

PTX 与 SASS

PTX：虚拟 ISA，类似于汇编语言，具有较好的可移植性
SASS：真实 GPU 指令集，针对特定架构优化

nvcc 默认生成 PTX 代码，在运行时由 GPU 驱动程序即时编译为 SASS。也可以使用 -arch 选项直接生成特定架构的 SASS 代码。

常用编译选项

-arch=sm_XX：指定目标 GPU 架构（如 sm_75 对应 Turing 架构）
-o output：指定输出文件名
-lcudart：链接 CUDA 运行时库
-Xcompiler：传递选项给主机编译器
-use_fast_math：使用快速数学函数（牺牲精度换取速度）

2.7 错误处理

CUDA API 调用可能失败，常见原因包括：

设备不存在或不可用
内存分配失败
数据传输错误
内核启动参数无效

检查 CUDA 错误

cpp

cudaError_t err = cudaGetLastError();
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
}

同步与错误检测

内核启动是异步的，错误可能在之后才显现。使用 cudaDeviceSynchronize() 等待所有操作完成后再检查错误：

cpp

kernel<<<grid, block>>>(args);
cudaError_t err = cudaDeviceSynchronize();
if (err != cudaSuccess) {
    printf("CUDA error: %s\n", cudaGetErrorString(err));
}

错误处理宏：

为了简化错误检查，可以定义宏：

cpp

#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error at %s:%d: %s\n", \
                   __FILE__, __LINE__, cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

使用：

cpp

CUDA_CHECK(cudaMalloc((void **)&d_data, size));
CUDA_CHECK(cudaMemcpy(d_data, h_data, size, cudaMemcpyHostToDevice));

三、程序员的视角

3.1 在实际项目中的应用

虽然"Hello World"本身没有实际应用价值，但它所展示的模式是所有 CUDA 程序的基础：

数据并行处理：图像处理、信号处理、科学计算等都遵循相同的数据传输-计算-返回模式
批量计算：金融模拟、机器学习推理等任务可以将数据分批发送到 GPU 处理
实时渲染：游戏引擎中的粒子系统、物理模拟等使用类似的内核启动模式

理解这一基本模式后，你可以将其应用到任何需要并行加速的场景中。

3.2 常见陷阱与误区

陷阱一：忘记检查错误

CUDA API 调用失败时不会抛出异常，而是返回错误码。如果不检查错误，程序可能静默失败，导致难以调试的问题。

陷阱二：内存泄漏

与 malloc/free 一样，cudaMalloc/cudaFree 必须配对使用。忘记释放设备内存会导致内存泄漏。

陷阱三：错误的线程索引计算

如果线程索引计算错误，可能导致数组越界访问或某些元素未被处理。始终确保边界检查：

cpp

int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) {
    // 安全访问
}

陷阱四：忽视异步特性

内核启动是异步的，这意味着主机代码在内核完成之前就会继续执行。如果需要确保内核完成，必须使用同步操作。

陷阱五：过度分配线程

启动比必要更多的线程不会提高性能，反而可能浪费资源。应根据数据规模合理配置线程数量。

3.3 与其他层级的联系

本章介绍的概念是后续章节的基础：

第3章：将深入探讨线程模型和执行配置
第4章：将详细介绍各种内存类型和管理策略
第5章：将讲解线程同步机制
第7章：将展示如何使用库来简化常见任务

四、代码实战

4.1 经典的 GPU Hello World

cpp

#include <stdio.h>

// 内核函数：每个线程打印一条消息
__global__ void helloFromGPU() {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    printf("Hello from thread %d!\n", tid);
}

int main() {
    printf("Hello from CPU!\n");
    
    // 启动 10 个线程
    helloFromGPU<<<1, 10>>>();
    
    // 同步以确保所有 printf 完成
    cudaDeviceSynchronize();
    
    return 0;
}

编译并运行：

bash

nvcc -o hello hello.cu
./hello

注意：GPU 上的 printf 输出可能会以任意顺序出现，因为线程是并行执行的。

4.2 向量加法完整示例

cpp

#include <stdio.h>
#include <stdlib.h>

// 向量加法内核
__global__ void vectorAdd(const float *a, const float *b, float *c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

// 验证结果
bool verifyResult(const float *c, const float *a, const float *b, int n) {
    for (int i = 0; i < n; i++) {
        if (fabs(c[i] - (a[i] + b[i])) > 1e-5) {
            return false;
        }
    }
    return true;
}

int main() {
    int n = 1 << 20; // 1M 元素
    size_t size = n * sizeof(float);
    
    // 分配主机内存
    float *h_a = (float *)malloc(size);
    float *h_b = (float *)malloc(size);
    float *h_c = (float *)malloc(size);
    
    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_a[i] = (float)i;
        h_b[i] = (float)i * 2.0f;
    }
    
    // 分配设备内存
    float *d_a, *d_b, *d_c;
    cudaError_t err;
    
    err = cudaMalloc((void **)&d_a, size);
    if (err != cudaSuccess) {
        printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    err = cudaMalloc((void **)&d_b, size);
    if (err != cudaSuccess) {
        printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    err = cudaMalloc((void **)&d_c, size);
    if (err != cudaSuccess) {
        printf("cudaMalloc failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    // 传输数据到设备
    err = cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        printf("cudaMemcpy failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    err = cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
        printf("cudaMemcpy failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    // 配置执行参数
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    
    // 启动内核
    vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
    
    // 检查内核执行错误
    err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        printf("Kernel execution failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    // 传输结果回主机
    err = cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
    if (err != cudaSuccess) {
        printf("cudaMemcpy failed: %s\n", cudaGetErrorString(err));
        return 1;
    }
    
    // 验证结果
    if (verifyResult(h_c, h_a, h_b, n)) {
        printf("Result verified successfully!\n");
    } else {
        printf("Result verification failed!\n");
    }
    
    // 清理
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

这个示例展示了完整的 CUDA 编程模式，包括错误处理、结果验证和资源清理。

4.3 使用 dim3 的多维线程配置

cpp

#include <stdio.h>

// 处理二维数据的内核
__global__ void process2DArray(float *data, int width, int height) {
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    
    if (col < width && row < height) {
        int idx = row * width + col;
        data[idx] = (float)(row * width + col);
    }
}

int main() {
    int width = 64;
    int height = 64;
    size_t size = width * height * sizeof(float);
    
    float *h_data = (float *)malloc(size);
    float *d_data;
    
    cudaMalloc((void **)&d_data, size);
    
    // 二维线程配置
    dim3 blockSize(16, 16);
    dim3 gridSize((width + blockSize.x - 1) / blockSize.x,
                  (height + blockSize.y - 1) / blockSize.y);
    
    process2DArray<<<gridSize, blockSize>>>(d_data, width, height);
    
    cudaMemcpy(h_data, d_data, size, cudaMemcpyDeviceToHost);
    
    // 打印部分结果
    for (int row = 0; row < 4; row++) {
        for (int col = 0; col < 4; col++) {
            printf("%.0f ", h_data[row * width + col]);
        }
        printf("\n");
    }
    
    free(h_data);
    cudaFree(d_data);
    
    return 0;
}

这个示例展示了如何使用多维线程配置来处理二维数据，这在图像处理和矩阵运算中非常常见。

4.4 错误处理宏

cpp

#include <stdio.h>

// 定义错误检查宏
#define CUDA_CHECK(call) \
    do { \
        cudaError_t err = call; \
        if (err != cudaSuccess) { \
            printf("CUDA error at %s:%d - %s\n", \
                   __FILE__, __LINE__, cudaGetErrorString(err)); \
            exit(1); \
        } \
    } while(0)

__global__ void simpleKernel() {
    int tid = threadIdx.x;
    printf("Thread %d running\n", tid);
}

int main() {
    float *d_data;
    CUDA_CHECK(cudaMalloc((void **)&d_data, 1024));
    
    simpleKernel<<<1, 8>>>();
    CUDA_CHECK(cudaDeviceSynchronize());
    
    CUDA_CHECK(cudaFree(d_data));
    
    printf("All operations completed successfully!\n");
    return 0;
}

使用宏可以简化错误检查代码，使主逻辑更加清晰。

五、深入思考

5.1 历史背景

CUDA 的编程模型受到了早期并行编程范式的深刻影响。SIMD 向量处理、数据并行语言（如 HPF）、以及图形着色器编程都为 CUDA 的设计提供了灵感。

NVIDIA 在设计 CUDA 时面临一个关键抉择：是提供一个接近硬件的低级 API，还是提供一个高度抽象的高级 API？最终的选择是折中方案：提供足够的抽象以简化编程，同时保留足够的底层控制以实现高性能。

5.2 前沿方向

当前的研究正在探索更高级的编程抽象：

统一内存：自动管理主机和设备之间的数据传输
C++ 标准支持：更好地集成现代 C++ 特性
图 API：将多个内核和数据传输操作组合为单个提交
异步编程模型：更好地利用 GPU 的并发能力

5.3 延伸阅读

NVIDIA CUDA C Programming Guide - 官方文档是最权威的资源
"CUDA by Example" by Jason Sanders and Edward Kandrot
CUDA Toolkit Documentation - 包含详细的 API 参考

六、实验笔记

实验 1：探索不同的线程配置

目标：观察不同线程块大小对性能的影响。

步骤：

使用向量加法示例

分别测试线程块大小为 32、64、128、256、512、1024

记录每种配置的执行时间

分析哪种配置性能最佳，并思考原因

提示：使用 cudaEvent 进行精确计时。

实验 2：测试边界条件

目标：验证当数据规模不是线程块大小的整数倍时的行为。

步骤：

修改向量加法示例中的数据规模为非 256 的倍数（如 1000003）

确认所有元素都被正确处理

尝试移除边界检查 if (idx < n)，观察会发生什么

思考题：为什么边界检查是必要的？在什么情况下可以省略？

七、本章小结

本章介绍了 CUDA 编程的基础知识。我们学习了主机与设备的分离模型，理解了内核函数的定义和启动方式，掌握了线程层次结构的概念，并熟悉了基本的内存管理 API。

关键的要点包括：

CUDA 程序遵循数据传输-计算-返回的基本模式
内核函数使用 __global__ 修饰符声明，通过 <<<grid, block>>> 语法启动
线程索引计算公式 idx = blockIdx.x * blockDim.x + threadIdx.x 是访问数据的基础
始终检查 CUDA API 调用的返回值以捕获错误
线程块大小的选择会影响性能，需要根据具体问题和硬件特性进行调整

通过动手实践本章的代码示例，你应该已经能够编写简单的 CUDA 程序。在下一章中，我们将深入探讨线程模型和执行配置的更多细节。

八、习题

解释主机和设备在 CUDA 编程模型中的角色分工。

写出启动一个包含 1024 个线程、每个线程块 256 个线程的内核函数的代码。

为什么内核函数的返回类型必须是 void？

描述 cudaMemcpy 的四种传输方向及其使用场景。

如果一个数组有 1000 个元素，线程块大小为 256，需要多少个线程块？最后一个线程块中有多少个活动线程？

编写一个 CUDA 内核，将数组中的每个元素乘以 2。

解释为什么内核启动是异步的，以及这种设计的优势。

nvcc 编译器如何处理主机代码和设备代码？

什么是 PTX？它与 SASS 有什么区别？

编写一个宏来简化 CUDA 错误检查，并在一个简单的程序中演示其用法。