3. CUDA 快速入门（Motivation: CUDA Crash Course）

3.1 GPU 与 CPU 的函数标识符

---

位置	函数修饰符	说明
GPU	__global__	全局核函数（kernel），可由 CPU 调用执行在 GPU 上）
GPU	__device__	设备函数，只能在 GPU 上被 GPU 调用
CPU	__host__	主机函数（默认是 CPU）

3.2 CUDA “Hello World” 示例（没有输出世界）

#include <cstdio>
// GPU 设备函数
__device__ int print() { return 0; }
// GPU 全局核函数
__global__ void kernel() { printf("%i", print()); }
// CPU 主机函数
__host__ void start() { kernel<<<2, 3>>>(); }
int main() { // 默认 __host__
    start();
    return cudaDeviceSynchronize();
}

理解：

1. __device__ int print()
   • 定义在 GPU 上的设备函数，返回 0。
2. __global__ void kernel()
   • GPU 核函数，可由 CPU 调用。
   • 内部调用 GPU 设备函数 print() 并通过 printf 打印返回值。
3. kernel<<<2, 3>>>();
   • CUDA 的执行配置：
     • 2 个块 (blocks)
     • 每个块 3 个线程 (threads per block)
   • CUDA 中核函数执行是并行的，线程编号可用 threadIdx、blockIdx 等标识。
4. cudaDeviceSynchronize()
   • 等待 GPU 所有核函数完成。
   • 返回 0 表示成功。
5. 编译器：
   • nvcc + 主机编译器（clang 或 gcc）可同时编译 CPU/GPU 代码。
   • 其他编译器如 HIP/nvc/gpucc 等也支持类似 CUDA 代码。

输出结果

stdout: 000000
return code: 0

解释：

• 由于 kernel 被 2 块 × 3 线程调用，每个线程都会打印 0，所以总共输出 6 个 0。
• return code: 0 表示程序正常退出，无错误。

3.3 CUDA 核函数的数学模型

假设有 $B$ 个块，每块 $T$ 个线程，总线程数为 $N=B\times T$。

• 每个线程独立执行 kernel：
  $${\text{Thread}}_{i,j}\text{ executes }f(x)i\in [0,B-1],j\in [0,T-1]$$
• 在本例中：
  $$B=2,T=3⟹N=6$$
  每个线程调用 print() 输出 0，所以总输出为 0 × 6。

3.4 关键注意点

1. CPU 代码通过 __host__ 调用 GPU 核函数 __global__。
2. GPU 内部函数只能在 GPU 上调用。
3. 并行执行会产生多个输出，需要注意 printf 的顺序可能不是固定的。
4. cudaDeviceSynchronize() 是 CPU 等待 GPU 完成的重要函数。

WSL（Windows Subsystem for Linux）下配置 CUDA 环境 的方法，并解释关键原理和注意事项。

1. WSL 支持的 CUDA 环境

• 要求：
  • Windows 10/11 的 WSL 2。
  • 安装 NVIDIA GPU 驱动 支持 WSL（通常是 515 及以上）。
  • WSL 内部 Linux 发行版（Ubuntu 推荐）通过 apt 安装 CUDA 工具包。
• CUDA 支持情况：
  • CUDA 在 WSL 下可以访问 GPU（类似 Linux 原生环境）。
  • 支持 nvcc 编译器、cuda-runtime、cuda-driver 等完整 CUDA 生态。
  • 可直接运行 GPU 核函数。

2. 安装步骤（以 Ubuntu 为例）

2.1 安装 WSL 2

wsl --install -d ubuntu
wsl --set-version ubuntu 2

• wsl --list --verbose 可以查看 WSL 版本。
• WSL 2 必须，因为 WSL 1 不支持 GPU。

2.2 安装 NVIDIA 驱动（Windows 端）

• 下载 支持 WSL 的 NVIDIA 驱动：
  • NVIDIA 官方 WSL 驱动下载
• 安装完成后，在 Windows CMD 中：

nvidia-smi

• 输出 GPU 信息表示驱动安装成功。

2.3 安装 CUDA 工具包（WSL 内 Ubuntu）

2.3.1 添加 NVIDIA 软件源

sudo apt update
sudo apt install -y wget gnupg
wget https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb
sudo dpkg -i cuda-keyring_1.1-1_all.deb
sudo apt update

2.3.2 安装 CUDA

sudo apt install -y cuda

会安装半天很多库下载

• 安装完成后，CUDA 默认在 /usr/local/cuda。
• 可用 nvcc --version 检查版本。

xiaqiu@xz:~$ whereis cuda
cuda: /usr/local/cuda
xiaqiu@xz:~$
xiaqiu@xz:~$ ls /usr/local/cuda/bin/nvcc
/usr/local/cuda/bin/nvcc
xiaqiu@xz:~$ nvcc
Command 'nvcc' not found, but can be installed with:
sudo apt install nvidia-cuda-toolkit

添加到环境变量

xiaqiu@xz:~$ echo 'export PATH=/usr/local/cuda/bin:$PATH' >> ~/.bashrc
xiaqiu@xz:~$ echo 'export LD_LIBRARY_PATH=/usr/local/cuda/lib64:$LD_LIBRARY_PATH' >> ~/.bashrc
xiaqiu@xz:~$ source ~/.bashrc
xiaqiu@xz:~$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Aug_20_01:58:59_PM_PDT_2025
Cuda compilation tools, release 13.0, V13.0.88
Build cuda_13.0.r13.0/compiler.36424714_0
xiaqiu@xz:~$

xiaqiu@xz:~$ ls /usr/local/cuda-13.0/bin/
__nvcc_device_query  cuda-gdb-minimal         cuda-gdbserver  nsight-sys                   nvdisasm
bin2c                cuda-gdb-python3.10-tui  cudafe++        nsight_ee_plugins_manage.sh  nvlink
compute-sanitizer    cuda-gdb-python3.11-tui  cuobjdump       nsys                         nvprune
crt                  cuda-gdb-python3.12-tui  fatbinary       nsys-ui                      ptxas
cu++filt             cuda-gdb-python3.8-tui   ncu             nvcc
cuda-gdb             cuda-gdb-python3.9-tui   ncu-ui          nvcc.profile
xiaqiu@xz:~$

在实际 CUDA 开发中，常用的工具其实集中在几个核心类别：编译、调试、性能分析和小工具。按用途划分如下：

1⃣ 编译相关（必用）

---

工具	用途
nvcc	核心编译器，把 .cu 文件编译成 GPU 可执行文件或库。
ptxas	PTX 汇编器，nvcc 内部使用，必要时可以单独查看 PTX 汇编生成 machine code。
nvlink	链接器，用于 device code 链接生成最终可执行文件或库。

开发中一般只直接用 nvcc，其他工具 nvcc 会自动调用。

2⃣ 调试相关（重要）

---

工具	用途
cuda-gdb	GPU kernel 调试器，单步调试 GPU 代码。
compute-sanitizer	检测内存越界、未初始化访问、race condition 等。
__nvcc_device_query	测试 GPU 是否可用，获取设备信息。

开发阶段必备：先用 device query 确认 GPU，逻辑复杂或报错用 compute-sanitizer，定位 bug 用 cuda-gdb。

3⃣ 性能分析（常用）

---

工具	用途
ncu	Nsight Compute 命令行版本，采集 GPU kernel 性能数据。
ncu-ui	Nsight Compute GUI，可视化分析 GPU kernel 性能。
nsys / nsys-ui	Nsight Systems（命令行/GUI），系统级性能分析，包括 CPU/GPU/IO。

性能优化阶段使用，普通开发可先用 ncu 或 ncu-ui。

4⃣ 二进制查看/小工具（偶尔用）

---

工具	用途
nvdisasm	查看 SASS / PTX 汇编，调试 GPU 汇编级问题。
cuobjdump	查看 GPU 对象文件，定位 PTX/SASS 信息。
cu++filt	还原 C++ mangled 名称，调试输出可读性。

一般普通开发很少用，除非调汇编或内核低层问题。

总结：常用工具 TOP 列表

1. 编译：nvcc
2. 调试：cuda-gdb, compute-sanitizer, __nvcc_device_query
3. 性能分析：ncu, ncu-ui
4. 偶尔：nvdisasm, cuobjdump

简单记忆：编译 → 调试 → 性能 → 汇编。

2.4 测试 CUDA

在 WSL 终端中创建 test.cu：

#include <cstdio>
__global__ void kernel() {
    printf("Hello CUDA from WSL!\n");
}
int main() {
    kernel<<<1,1>>>();
    cudaDeviceSynchronize();
    return 0;
}

编译并运行：

nvcc test.cu -o test
./test

# CMakeLists.txt — 现代方式（推荐）
cmake_minimum_required(VERSION 3.18) # 3.18+ 更好，3.10 也可
# CUDA 架构：可设置为 AUTO 或指定 compute capability 列表，例如 70;75;80
if(NOT DEFINED CUDA_ARCHS)
  set(CUDA_ARCHS native)  
endif()
set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCHS})
project(MyCudaProject LANGUAGES CXX CUDA)
# C/C++ 标准
set(CMAKE_CXX_STANDARD 20)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_STANDARD 14)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
# 编译选项（示例）
add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-fPIC>")
# 查找源码（修改成你自己的源文件）
set(SRC
    main.cu
)
# 可选：把 CUDA separable compilation 打开（如果有 device 函数跨文件调用）
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
# 可选：如果要生成位置无关代码（对于 shared lib 有用）
set(CMAKE_POSITION_INDEPENDENT_CODE ON)
# 创建可执行文件 / 库
add_executable(my_app ${SRC})
# 如果你的 CUDA 运行时需要链接（通常 CMake 会自动处理）
find_package(CUDAToolkit REQUIRED) # modern CMake module
target_link_libraries(my_app PRIVATE CUDA::cudart)
# 包含目录
target_include_directories(my_app PRIVATE ${CMAKE_SOURCE_DIR}/include)
# 目标属性（示例）
target_compile_features(my_app PRIVATE cxx_std_20)
# 如果你想让 nvcc 使用特定编译选项：
target_compile_options(my_app PRIVATE
  "$<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>"
)
# 安装规则（可选）
install(TARGETS my_app RUNTIME DESTINATION bin)

• 如果输出 Hello CUDA from WSL!，说明环境配置成功。

3. WSL CUDA 的注意事项

1. 驱动版本匹配：
   • Windows 端 GPU 驱动必须支持 WSL。
   • WSL 内的 CUDA 版本需要和驱动兼容。
2. 文件系统性能：
   • WSL 读取 Windows 文件 (/mnt/c/...) 比较慢，建议代码在 WSL 内 /home/... 下。
3. GUI/OpenGL：
   • CUDA 本身可以用，但图形渲染或 GUI 可能需要 WSLg 或额外配置。
4. 多 GPU 支持：
   • 可用 nvidia-smi 或 cudaGetDeviceCount() 检查 GPU 数量。

4. 数学/原理理解

WSL 下 CUDA 原理与 Linux 类似：

• CPU（主机）在 WSL 内发起核函数调用：
  $$\text{kernel}<<<B,T>>>()$$
• 通过 Windows 驱动与 GPU 通信，GPU 执行线程：
  $$N=B\times T\text{线程并行计算}$$
• 数据在 WSL 内存空间 和 GPU 显存之间传输：
  $$\text{Host memory}\stackrel{\text{cudaMemcpy}}{\leftrightarrow }\text{Device memory}$$

发现代码没有打印输出"Hello CUDA from WSL!"

#include <cstdio>
#include <cuda_runtime.h>
__global__ void kernel() {  //
    std::printf("Hello CUDA from WSL!\n");
}
int main() {
    kernel<<<1, 1>>>();
    cudaError_t err = cudaDeviceSynchronize();
    if (err != cudaSuccess) {
        printf("CUDA Error: %s\n", cudaGetErrorString(err));
    }
    return 0;
}

CUDA driver version 和 CUDA runtime version 不匹配

CUDA Error: CUDA driver version is insufficient for CUDA runtime version

xiaqiu@xz:~/test$ nvidia-smi
Wed Dec  3 14:32:30 2025       
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 580.105.08             Driver Version: 566.07         CUDA Version: 12.7     |
+-----------------------------------------+------------------------+----------------------+
| GPU  Name                 Persistence-M | Bus-Id          Disp.A | Volatile Uncorr. ECC |
| Fan  Temp   Perf          Pwr:Usage/Cap |           Memory-Usage | GPU-Util  Compute M. |
|                                         |                        |               MIG M. |
|=========================================+========================+======================|
|   0  NVIDIA GeForce RTX 4060 ...    On  |   00000000:01:00.0 Off |                  N/A |
| N/A   41C    P8              3W /   95W |     725MiB /   8188MiB |      2%      Default |
|                                         |                        |                  N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes:                                                                              |
|  GPU   GI   CI              PID   Type   Process name                        GPU Memory |
|        ID   ID                                                               Usage      |
|=========================================================================================|
|  No running processes found                                                             |
+-----------------------------------------------------------------------------------------+
xiaqiu@xz:~/test$ 
xiaqiu@xz:~/test$ nvcc --version
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Wed_Aug_20_01:58:59_PM_PDT_2025
Cuda compilation tools, release 13.0, V13.0.88
Build cuda_13.0.r13.0/compiler.36424714_0
xiaqiu@xz:~/test$ 

CUDA Version: 12.7
nvcc --version release 13.0, V13.0.88

如何更新 Windows 驱动（步骤详解）

问题找到了！CUDA 运行时版本（13.0）比驱动支持的版本（12.7）要新。
驱动支持：CUDA 12.7
运行时版本：CUDA 13.0
这就是不匹配的原因。

解决方案 ：更新 Windows 驱动（如果可行）

在 Windows 中更新到支持 CUDA 13.0 的最新驱动：
最好方法windows的nvidia app更新最新的显卡驱动

1. 重启 WSL：

wsl --shutdown

然后重新打开 WSL

临时解决方案：使用已有的旧版本 CUDA（如果存在）

检查是否有其他 CUDA 版本：

ls /usr/local/ | grep cuda

如果有 cuda-12.x，可以临时切换：

export PATH=/usr/local/cuda-12.6/bin:$PATH
export LD_LIBRARY_PATH=/usr/local/cuda-12.6/lib64:$LD_LIBRARY_PATH
nvcc --version  # 验证版本

我的方法是用nvidia app 更新最新的显卡驱动

C:\Users\16956>nvidia-smi
Wed Dec  3 15:27:29 2025
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 581.57                 Driver Version: 581.57         CUDA Version: 13.0     |
+-----------------------------------------+------------------------+----------------------+
| GPU  Name                  Driver-Model | Bus-Id          Disp.A | Volatile Uncorr. ECC |
| Fan  Temp   Perf          Pwr:Usage/Cap |           Memory-Usage | GPU-Util  Compute M. |
|                                         |                        |               MIG M. |
|=========================================+========================+======================|
|   0  NVIDIA GeForce RTX 4060 ...  WDDM  |   00000000:01:00.0 Off |                  N/A |
| N/A   45C    P0             18W /   95W |       0MiB /   8188MiB |      0%      Default |
|                                         |                        |                  N/A |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes:                                                                              |
|  GPU   GI   CI              PID   Type   Process name                        GPU Memory |
|        ID   ID                                                               Usage      |
|=========================================================================================|
|  No running processes found                                                             |
+-----------------------------------------------------------------------------------------+
C:\Users\16956>

再次运行程序输出

Hello CUDA from WSL!

windows 下的Nsight Compute可以下载安装

1⃣ 下载 Nsight Compute

1. 打开 NVIDIA 官方下载页：NVIDIA Nsight Compute
2. 选择 Nsight Compute 版本
   • Windows 版：下载 .exe 安装包
   • Linux 版（如果在 WSL 或 Linux 使用）：下载 .run 或 .deb 文件

注意：Nsight Compute 版本要与你的 CUDA Toolkit 驱动兼容，最好使用最新版本或与 CUDA 12.x/13.x 对应的版本。

2⃣ Windows 安装步骤

1. 双击下载的 NsightCompute-win-x.y.z.exe 安装包
2. 选择安装目录（默认即可）
3. 安装完成后，会生成以下主要内容：
   • GUI 应用：用于性能分析和查看报告
   • 命令行工具 ncu：可在命令行使用
   • 集成 Visual Studio 插件（可选）
4. 配置环境变量（可选）：
   • 将 Nsight Compute 的 bin 目录加入 PATH，例如：

     C:\Program Files\NVIDIA Corporation\Nsight Compute 2025.1\bin
     

   • 这样就可以在 CMD / PowerShell 里直接运行 ncu.exe

**3⃣ Linux / WSL **

xiaqiu@xz:~$ whereis ncu-ui
ncu-ui: /usr/local/cuda-13.0/bin/ncu-ui
xiaqiu@xz:~$
xiaqiu@xz:~$ ncu-ui
xiaqiu@xz:~$

在 WSLg（Windows Subsystem for Linux GUI） 下，ncu-ui 的界面缩小或者字体很小，是常见问题，因为 WSLg 的 DPI 缩放和 Qt 应用的默认缩放有关。你可以通过以下方法放大界面或字体：

1⃣ 设置环境变量 QT_SCALE_FACTOR

ncu-ui 是 Qt 应用，可以通过 Qt 的缩放变量放大：

export QT_SCALE_FACTOR=1.5   # 1.5 倍缩放，可以改成 2
ncu-ui

• 1.0 → 默认大小
• 1.5 → 放大 1.5 倍
• 2.0 → 放大 2 倍

可以在 ~/.bashrc 或 ~/.zshrc 添加这一行，让每次打开自动生效。

2⃣ 使用 QT_AUTO_SCREEN_SCALE_FACTOR

Qt 可以自动根据 DPI 缩放：

export QT_AUTO_SCREEN_SCALE_FACTOR=1
ncu-ui

• 会根据 WSLg 提供的屏幕 DPI 自动放大
• 有时比 QT_SCALE_FACTOR 更智能

3⃣ 改变 Windows 缩放设置

如果你在 4K 显示器下，Windows 的显示缩放会影响 WSLg 窗口：

1. 右键桌面 → 显示设置 → 缩放与布局
2. 将缩放比例调整到 125%~200%，然后重新打开 ncu-ui

4⃣ 临时命令一行启动

结合以上方法，可以直接在终端运行：

QT_SCALE_FACTOR=1.8 ncu-ui

建议：

• 如果你经常用 ncu-ui，最好把 QT_SCALE_FACTOR=1.5 加到 ~/.bashrc，这样每次打开都会按你设定的缩放显示。
• 如果界面仍然很小，配合 WSLg DPI 设置 + QT_SCALE_FACTOR 一般就能正常。

1⃣ 代码分析

#include <cstdio>                     // 1
__device__ int print() { return 0; }  // 3
__global__ void kernel() {            // 4
    std::printf("%i", print());
}
__host__ void start() {  // 5
    kernel<<<2, 3>>>();  // 启动 kernel，grid=2, block=3
}
int main() {                         // 7
    start();                         // 8
    return cudaDeviceSynchronize();  // 9-10
}

关键点解释

1. __device__ 函数

__device__ int print() { return 0; }

• 这是 设备端函数，只能在 GPU kernel 或其他 device 函数里调用。
• 返回值是 0。
• 数学理解：
  • 每个线程执行 print() 得到的值是：
    $$f_{\text{device}}(threadIdx,blockIdx)=0$$

2. __global__ kernel

__global__ void kernel() { printf("%i", print()); }

• __global__ 表示这是 GPU kernel，可以被 host 调用。
• 内部调用 print() 并打印结果。
• 注意：printf 在 GPU 内是异步输出的，实际输出需要 cudaDeviceSynchronize() 才会刷新。

3. 启动 kernel

kernel<<<2, 3>>>();

• <<<gridDim, blockDim>>> 是 CUDA 启动 kernel 的语法。
• gridDim=2 → 两个 block
• blockDim=3 → 每个 block 三个线程
• 总线程数：
  $$N_{\text{threads}}=\text{gridDim}\times \text{blockDim}=2\times 3=6$$
• 每个线程都会调用一次 print() 并打印 0，因此输出：

000000

4. __host__ 函数 start

__host__ void start() { kernel<<<2,3>>>(); }

• host 端函数，用来启动 kernel。
• CUDA kernel 是 异步执行，调用 kernel 后会立即返回，kernel 在 GPU 上运行。

5. 同步

return cudaDeviceSynchronize();

• cudaDeviceSynchronize() 等待 GPU 上所有 kernel 执行完成，并收集错误码。
• 这里保证所有线程的 printf 输出在 host 端可见。

2⃣ 程序执行流程（理解）

1. 主函数 main() 调用 start()
2. start() 启动 kernel：
   • grid=2，block=3 → 总 6 个线程
   • 每个线程调用 print() → 输出 0
3. GPU kernel 异步执行
4. 主函数调用 cudaDeviceSynchronize() → 等待 kernel 执行完成
5. 输出：

000000

• 每个 0 对应一个线程输出
• 返回值 0 表示程序成功结束

3⃣ 数学/线程视角理解

• 总线程数 $N_{\text{threads}}$：
  $$N_{\text{threads}}=\text{gridDim}\times \text{blockDim}=2\times 3=6$$
• 每个线程的执行过程：
  $$\text{thread }t:x_t=f_{\text{device}}(t)=0$$
• 最终输出可以看作：
  $$\text{Output}=\prod _{t=1}^6{x}_t$$
  在这里，x_t = 0，打印显示为 000000。

4⃣ 关于输出和调试

• printf 在 kernel 内异步输出
• 如果没有 cudaDeviceSynchronize()，可能看不到输出
• 调试 kernel 时：
  • 不能直接在 __device__ 函数打断点
  • 可以在 kernel 内设置断点
• 输出顺序可能不是严格按照线程索引，因为 GPU 线程是并行执行的

5⃣ 总结

1. CUDA kernel 是 GPU 并行执行单元，每个线程独立运行
2. __device__ 函数只能在 GPU 上调用
3. kernel<<<grid, block>>> 启动 kernel，grid 和 block 决定总线程数
4. cudaDeviceSynchronize() 保证 kernel 执行完成和 printf 输出可见
5. 本例输出：

000000

1⃣ 代码回顾

// 定义一个只在 Host（CPU）上可调用的结构体 H
struct H {
    __host__ int func() { 
        return 42;               // Host 函数 func，返回整数 42
    } 
};  // 1
// 定义一个只在 Device（GPU）上可调用的结构体 D
struct D {
    __device__ int func() { 
        return 666;              // Device 函数 func，返回整数 666
    } 
};  // 3
// 定义一个模板 wrap，可以同时在 Host 和 Device 上调用
template <typename T>
__host__ __device__ int wrap() {
    // 调用类型 T 的 func 方法
    // 注意：
    // 1. 如果 T=H → host 版本 ok，device 版本非法（警告）
    // 2. 如果 T=D → host 版本非法（编译错误或 UB），device 版本 ok
    return T{}.func();
}  // 5-6
int main() {
    // Host 调用 H 的 func → 合法
    // return H{}.func();  // OK
    // Host 调用 D 的 func → 非法
    // return D{}.func();  // 编译错误：host 不能调用 device 函数
    // Host 调用 wrap<H>() → 会生成 device 版本 wrap<H>()，对 device 来说 H.func() 非法 → 编译警告
    // return wrap<H>();   // 警告
    // Host 调用 wrap<D>() → host 版本 wrap<D>() 内调用 D.func() 非法 → UB / 运行时未定义行为
    // return wrap<D>();   // UB / 运行时未定义行为
}

2⃣ CUDA 的 host/device 函数规则

1. __host__ 函数
   • 只能在 CPU 端 调用
   • 对应主机执行环境
2. __device__ 函数
   • 只能在 GPU 端 调用
   • 对应 device 执行环境
   • 不能直接在 CPU 端调用
3. __host__ __device__ 函数
   • 可以同时在 CPU 和 GPU 调用
   • 编译器会生成两份函数版本：
     1. host 版本（CPU 调用）
     2. device 版本（GPU 调用）
   • 注意：模板实例化时，调用的成员函数必须符合目标环境的规则

3⃣ wrap 模板函数解析

template<typename T> __host__ __device__
int wrap() { return T{}.func(); }

• 当 wrap<H>() 在 host 端 被调用：
  • 生成 host 版本：

    int wrap<H>() { return H{}.func(); }  // OK
    

  • 生成 device 版本：

    __device__ int wrap<H>() { return H{}.func(); } // 错误，device 不能调用 host
    

  • 所以会报 警告 #20014-D，因为 device 版本不能调用 host 函数
• 当 wrap<D>() 在 host 端 被调用：
  • host 版本：

    int wrap<D>() { return D{}.func(); } // 错误，host 不能调用 device
    

  • device 版本：

    __device__ int wrap<D>() { return D{}.func(); } // OK
    

  • 所以如果只调用 host 版本，会编译报错或 UB

4⃣ main 函数中不同调用行为

---

调用	说明	结果
H{}.func()	host 调用 host 函数	OK
D{}.func()	host 调用 device 函数	编译错误
wrap<H>()	host 调用 wrap()	警告，device 版本非法
wrap<D>()	host 调用 wrap()	UB/运行时未定义行为

5⃣ 数学公式角度理解

1. 函数调用限制可以看作集合关系：

• 定义集合：
  • $H_{\text{host}}$ = { host 函数 }
  • $D_{\text{device}}$ = { device 函数 }
• 允许调用规则：
  $$\text{host 环境只能调用 }{H}_{\text{host}}\text{device 环境只能调用 }{D}_{\text{device}}$$

2. 模板 wrap 生成两份版本：

• Host 版本：
  $$wra{p}_H^{\text{host}}=H_{\text{host}}\text{ (合法) }wra{p}_D^{\text{host}}=D_{\text{device}}\text{ (非法) }$$
• Device 版本：
  $$wra{p}_H^{\text{device}}=H_{\text{host}}\text{ (非法) }wra{p}_D^{\text{device}}=D_{\text{device}}\text{ (合法) }$$

这就是为什么 wrap<H>() 会产生警告，wrap<D>() 会在 host 调用时报错或 UB。

6⃣ 总结

1. Host/Device 函数不能跨环境调用
   • host 不能直接调用 device
   • device 不能直接调用 host
2. __host__ __device__ 模板
   • 编译器会生成两份函数版本
   • 调用模板时，必须确保调用环境和目标函数匹配
3. 常见错误模式：

D{}.func()   // host 端调用 device → 编译错误
wrap<D>()    // host 端实例化 device 函数 → UB

4. 编译器提示：

• 警告 #20014-D：调用 host 函数在 device 版本
• 错误 #20014-D：调用 device 函数在 host 版本

5. 编译器和平台差异：

• nvcc 12.6 遵守严格规则
• Clang / HIP 编译器对 host/device 检测略有不同，有些可能允许，但运行时行为未定义

1⃣ “Oldschool” 方法——旧版本 CUDA

在早期 CUDA 中，host/device 函数的使用有很多限制：

• 只能用很老的 CUDA 版本才能实现跨 host/device 的模板调用。
• host/device 关键字在普通 C++ 编译器（如 gcc）下会报错，因为 gcc 不认识 __host__ 和 __device__。

2⃣ host / device / all the things 模式

这个模式的核心是：

1. host → CPU 可调用函数
2. device → GPU 可调用函数
3. all the things → 同时在 CPU 和 GPU 可调用函数（__host__ __device__）
   示例：

__host__ __device__ inline constexpr
float MHzToWavelength(int frequency_in_MHz) {
    return 299792458.f / (frequency_in_MHz * 1000.f * 1000.f);
}

• 将函数标记为 __host__ __device__：
  • CPU 上可直接调用
  • GPU 上可在 kernel 内调用
• constexpr 可以在编译期计算，适合常量计算
  数学上：
  $$\lambda =\frac{c}{f}=\frac{299792458}{f_{\text{Hz}}}\text{其中 }{f}_{\text{Hz}}=f_{\text{MHz}}\cdot 10^6$$

3⃣ 跨平台编译问题

如果直接用 gcc 编译上面的函数：

gcc mhz.cpp

• 会报错：

error: ‘__host__’ does not name a type

• 原因：gcc 不认识 __host__ 和 __device__，这些是 CUDA nvcc 扩展关键字。
  解决方法：条件宏定义

#ifndef CUDATAGS
#define CUDATAGS
#ifndef __CUDACC__  // 如果不是 nvcc 编译
#define __host__
#define __device__
#endif
#endif

• 这样在非 CUDA 编译器下，__host__ 和 __device__ 被定义为空，保证普通 C++ 编译器能编译。
• 使用示例：

__host__ __device__
void func() {}

• 在 GPU 上 nvcc 编译生成 device/host 两份版本
• 在 CPU 上普通编译器可直接调用

4⃣ Host / Device / Everywhere 示例

struct H {
    __host__ int func() { return 42; }
};
struct D {
    __device__ int func() { return 666; }
};
template<typename T>
int wrap() { return T{}.func(); }
int main() {
    H{}.func();      // OK
    D{}.func();      // GPU 端 OK, CPU 端不可
    wrap<H>();       // OK
    wrap<D>();       // Device 内 OK, Host 内 UB
}

核心概念：

• Host 调用 device 函数 → 编译错误
• Device 调用 host 函数 → 编译错误
• wrap<T> 模板必须确保类型 T 对应的函数在目标环境合法

5⃣ 条件函数体 / 条件编译

• 可以使用宏定义简化：

#ifndef __CUDACC__
#define HST
#define DEV
#else
#define HST __host__
#define DEV __device__
#endif
HST DEV void func() {}

• HST DEV 在 CPU 上为空，在 CUDA 上是 __host__ __device__
• 避免在非 CUDA 编译器报错

6⃣ 数学 / 物理公式示例

以 MHz 转波长函数为例：

1. 波长公式：
   $$\lambda =\frac{c}{f}$$
2. 将 MHz 转为 Hz：
   $$f_{\text{Hz}}=f_{\text{MHz}}\cdot 10^6$$
3. 综合公式：
   $$\lambda =\frac{299792458}{f_{\text{MHz}}\cdot 10^6}\text{单位：米 (m)}$$

• 使用 constexpr 可以在编译期计算波长，提高性能

7⃣ 总结

1. Host / Device / Everywhere 模式
   • __host__ __device__ → 生成 CPU/GPU 两份版本
   • 适合模板函数、物理计算函数
2. 条件编译
   • 使用宏判断是否 nvcc
   • 在非 CUDA 编译器下定义为空，保证可移植性
3. 注意事项
   • Host 不能调用 Device 函数
   • Device 不能调用 Host 函数
   • 模板 wrap() 要确保每个实例化版本合法，否则可能 UB
4. 数学公式
   • 将物理公式（如波长）直接写成 constexpr 函数，可同时在 CPU/GPU 上计算
   • $$\lambda =\frac{c}{f_{\text{Hz}}}=\frac{299792458}{f_{\text{MHz}}\cdot 10^6}$$

1⃣ 条件函数体的概念

在 CUDA 中，我们经常希望 同一函数在 Host（CPU）和 Device（GPU）端使用不同实现。
例如：计算向量的 2-范数：

struct Vec3f { float x, y, z; float operator[](int i) const; };

• Host 实现（在 CPU 上）：手动累加平方和再开方

float norm(const Vec3f &v) {
    double sum = 0;
    for (int i = 0; i < 3; ++i)
        sum += v[i]*v[i];
    return sqrtf((float)sum);
}

• Device 实现（在 GPU 上）：调用 CUDA 内置函数 norm3df

float norm(const Vec3f &v) {
    return norm3df(v[0], v[1], v[2]);
}

2⃣ 条件编译实现

为了写同一个函数名 norm()，可以使用宏判断 __CUDA_ARCH__：

__host__ __device__
float norm(const Vec3f &v) {
#ifndef __CUDA_ARCH__   // CPU 端
    double sum = 0;
    for (int i = 0; i < 3; ++i)
        sum += v[i]*v[i];
    return sqrtf((float)sum);
#else                   // GPU 端
    return norm3df(v[0], v[1], v[2]);
#endif
}

• 宏 __CUDA_ARCH__ 在 device 编译时定义
• 宏判断保证 Host / Device 使用不同函数体
• 公式解释：
  $$||\mathbf{v}|{|}_2=\sqrt{v_0^2+v_1^2+v_2^2}$$
  在 CPU 上手算，GPU 上调用 norm3df 函数（硬件加速）

3⃣ 条件函数体的问题

1. Host / Device 不同版本函数签名限制：
   • 函数签名、模板实例化参数 不能依赖 __CUDA_ARCH__
   • 如果用 if constexpr(__CUDA_ARCH__) 或宏修改签名，nvcc 会报错或 UB
   • 仅能修改函数体内容，不要修改返回类型或参数类型
2. 示例错误用法

struct H { __host__ int func() { return 42; } };
template<typename T>
__host__ __device__
void wrap() { 
    #ifndef __CUDA_ARCH__
    T{}.func();   // 在 CPU 编译时 OK
    #else
    T{}.func();   // Device 编译时可能 UB
    #endif
}

• 在 wrap 模板实例化时，如果 T 类型在 Device 上不可调用 host 函数 → UB
• 所以不推荐在模板中使用宏条件改变函数体调用

4⃣ 条件返回类型的安全做法

有时候希望 Host / Device 函数返回不同类型：

• 错误示例

__host__ H func() { return H{}; }
__device__ D func() { return D{}; }

• 在 nvcc 编译时，如果同时在模板或宏中调用可能产生 UB
• 推荐做法

#ifndef __CUDA_ARCH__
__host__ H func() { return H{}; }
#else
__device__ D func() { return D{}; }
#endif

• Host 端只看到 H 类型函数
• Device 端只看到 D 类型函数
• 保证类型安全和调用合法性

5⃣ 总结

1. Host / Device 条件函数体：
   • 使用 #ifndef __CUDA_ARCH__ / #else 判断函数体内容
   • 不要改变函数签名或模板参数
   • CPU 上手动实现，GPU 上调用内置函数（如 norm3df）
2. 模板函数安全使用：
   • 模板 wrap() 调用必须确保 T 的 func() 在目标环境合法
   • 宏条件只能影响函数体，不影响模板类型或返回类型
3. 条件返回类型：
   • 不要在同一函数名下返回不同类型
   • 用宏区分 host/device 函数定义即可

1⃣ 背景

在 C++ 中，constexpr 函数表示 可以在编译期求值，通常要求函数体内只调用 constexpr 函数或可编译期求值的操作。
在 CUDA 中，如果你想在 Device（GPU）端也使用 constexpr，就会遇到 nvcc 的限制：

constexpr int func() {
    std::array<int,5> a;
    // ...
    return std::accumulate(a.begin(), a.end(), 0);
}

• 编译时 nvcc 报错：

error: calling a host function("std::accumulate") from a device function("func") is not allowed

• 原因：std::accumulate 是 host 函数，Device 端不允许直接调用

2⃣ 解决方案：–expt-relaxed-constexpr

CUDA 提供了一个实验性编译选项：

--expt-relaxed-constexpr

作用：

1. 允许更多 host-side constexpr 函数在 device 上求值
2. 让一些原本不允许的 constexpr 调用在 nvcc 编译时通过
3. 不需要修改源代码，但属于实验性功能，可能有潜在风险
   例如：

nvcc --expt-relaxed-constexpr test.cu

这会允许像 std::array、std::accumulate 等在 device 上的 constexpr 调用。

3⃣ 已知使用案例

1. RAPIDS1（NVIDIA 数据科学和 AI 库）
   • 讨论是否使用 --expt-relaxed-constexpr，最终决定不使用
2. MatX2（NVIDIA 矩阵库）
   • 使用了 --expt-relaxed-constexpr
3. Dimetor
   • 仍在考虑使用

4⃣ 失败示例

• 在 nvcc 12.2 之前，如果在 constexpr 函数里调用非 constexpr 函数或抛异常：

constexpr int foo(int j) {
    if (j < 0) throw;   // 或 return bar(j);
}

• nvcc 会直接忽略原函数内容：

__device__ constexpr int foo(int j) { return 42; }

• 也就是说，函数行为被重写，可能产生 未定义行为（UB）
• 使用 malloc/free 也会失败：

constexpr int set() {
    auto i = (int*) malloc(sizeof(int));
    *i = 42;
    int y = *i;
    free(i);
    return y;
}

• 是否能工作取决于 编译器/操作系统/驱动/GPU
• 并不是编译期求值真正意义上的 constexpr

5⃣ 后果与注意事项

优点：

1. 易用：几乎不用改动代码
2. 支持第三方 constexpr 函数：std::array、std::optional、std::upper_bound 等

缺点：

1. 只适用于 constexpr 函数
2. 如果在库中使用可能导致不兼容或隐藏 bug
3. 属于实验性功能（<= 2016），未来 C++ 版本可能不支持
4. 可能引发 subtle bugs（微妙的运行时错误）

6⃣ 总结

• 问题：Device 端无法调用 host-side constexpr 函数
• 实验性解决方案：使用 nvcc --expt-relaxed-constexpr
• 使用场景：
  • 数值库、矩阵库、科学计算库中常见
  • 需要 minimal 代码修改，但要注意兼容性
• 风险：
  • 实验性特性 → 未来可能不稳定
  • 可能隐藏运行时问题

1⃣ 问题背景

你最开始写了一个模板函数：

template< typename Container >
__host__ __device__ constexpr
void fill_ones(Container & ct) {
    for (auto & x : ct) x = 1;
}

然后在 main 里用：

std::vector<int> v{1,2,3,4,5};
fill_ones(v);

结果：

warning #20014-D: calling a __host__ function from a __host__ __device__ function is not allowed

原因：

• __host__ __device__ 表示这个函数既可以在 CPU 上调用，也可以在 GPU 上调用。
• std::vector 的操作（如 begin()、operator[] 等）是 host-only 的。
• 因此，nvcc 会在 __device__ 上生成代码时报错（或至少警告）。
  换句话说，你不能在 device 函数里调用 host-only 函数。

2⃣ 禁用警告的几种方法

(a) #pragma 方法

#pragma nv_diag_suppress 20014

或者针对多个警告：

#pragma nv_diag_suppress 20011,20014

• 作用：告诉 nvcc 编译器不要显示某些警告。
• 位置：函数或文件开头。
• 缺点：
  • 这是 undocumented 的 pragma，未来可能会失效。
  • 不能解决根本问题，只是隐藏警告。

(b) nv_exec_check_disable（Thrust、Eigen 等常用）

#pragma nv_exec_check_disable

• 用于 Thrust 或 Eigen 的 host/device 函数模板中，忽略某些 exec-check。
• 也有类似 --expt-relaxed-constexpr 选项用于 Eigen，使 constexpr 与 device 函数兼容。

© 编译器全局选项

--diag-suppress 20011,20014

• 在 nvcc 命令行加上，可以禁掉指定警告。

3⃣ 根本解决方案：模板分离 Host/Device

警告的根源是 函数被错误实例化，调用了 host-only 的代码。
正确做法是生成不同版本的函数。

(a) 定义 Host/Device 兼容性类型

enum class HDC { Hst, Dev, HstDev }; // Host / Device / Both

(b) 为不同目标写不同版本函数

template<typename T, HDC x = hdc<T>>
requires(x == HDC::Hst)
__host__ void func(T t) { /* Host-only body */ }
template<typename T, HDC x = hdc<T>>
requires(x == HDC::Dev)
__device__ void func(T t) { /* Device-only body */ }
template<typename T, HDC x = hdc<T>>
requires(x == HDC::HstDev)
__host__ __device__ void func(T t) { /* Host+Device */ }

• hdc 是一个 trait，用来判断类型 T 的目标：
  • 例如 std::vector<int> 只能用于 Hst
  • float[N] 或自定义数组可用于 Dev 或 HstDev

© 包装调用函数

template<typename T>
__host__ __device__ void wrap(T t) {
    func(t);  // 根据 T 的类型实例化正确版本
}

这样：

• CPU 调用 wrap(v) → 调用 host 版本。
• GPU 调用 wrap(arr) → 调用 device 版本。
• 避免了 host-only 调用被实例化到 device 的情况。

4⃣ 数学/类型理解（用公式表示）

如果把类型目标映射抽象为数学公式，可以写为：
$${\text{func}}_{\text{target}}:T\mapsto \{\begin{array}{ll}\text{host-only},& T\in H\\ \text{device-only},& T\in D\\ \text{host+device},& T\in H\cap D\end{array}$$

• $H$ 表示 host-only 类型集合
• $D$ 表示 device-compatible 类型集合
• $H\cap D$ 表示两边都可用的类型

wrap 函数相当于：

$$\text{wrap}(T)={\text{func}}_{\text{hdc}(T)}(T)$$

5⃣ 使用示例

#include <vector>
std::vector<int> v{1,2,3,4,5};
// Host 调用
wrap(v);
// Device 调用
float arr[5] = {0,0,0,0,0};
wrap(arr);  // 生成 __device__ 版本

• 注意：在 kernel 中不能传 std::vector：

__global__ void kernel() {
    // fill_ones(H{}); // 错误：H 是 host-only
    fill_ones(D{});     // 正确：D 是 device-compatible
}

总结

1. 不要直接在 __host__ __device__ 函数里用 std::vector 或 host-only 类型。
2. 禁用警告可以用 #pragma nv_diag_suppress 或编译器选项，但只是表面解决。
3. 根本解决方案是：
   • 用模板区分 Host / Device / Both
   • 根据类型实例化对应版本
   • 用 wrapper 函数统一调用
     这样既 安全 又 可维护。

#include <iostream>
#include <vector>
#include <type_traits>
// ==========================
// Host/Device 类型 trait
// ==========================
enum class HDC { Hst, Dev, HstDev };
// 判断类型是否 host-only
template <typename T>
struct hdc {
    static constexpr HDC value = std::is_array_v<T> ? HDC::HstDev : HDC::Hst;
};
// ==========================
// fill_ones 函数
// ==========================
template <HDC tag, typename Container>
__host__ __device__ void fill_ones_impl(Container& ct) {
    if constexpr (tag == HDC::Hst || tag == HDC::HstDev) {
        // CPU 版本
        for (auto& x : ct) x = 1;
    }
    // 如果需要 GPU 版本，可以在这里添加
    // if constexpr (tag == HDC::Dev || tag == HDC::HstDev) {
    //     // GPU 版本实现
    // }
}
template <typename T, HDC tag = hdc<T>::value>
__host__ __device__ void fill_ones(T& ct) {
    fill_ones_impl<tag>(ct);
}
// ==========================
// CUDA kernel 示例
// ==========================
__global__ void kernel_fill(float* arr, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) arr[idx] = 1.0f;
}
// ==========================
// 主程序示例
// ==========================
int main() {
    // ===== CPU host-only 容器 =====
    std::vector<int> v{1, 2, 3, 4, 5};
    std::cout << "Before fill_ones (vector): ";
    for (auto x : v) std::cout << x << " ";
    std::cout << "\n";
    fill_ones(v);  // 安全调用 host-only vector
    std::cout << "After fill_ones (vector): ";
    for (auto x : v) std::cout << x << " ";
    std::cout << "\n";
    // ===== CPU/GPU 兼容容器 =====
    float arr[5] = {0, 0, 0, 0, 0};
    fill_ones(arr);  // 生成 host+device 版本
    std::cout << "After fill_ones (array): ";
    for (auto x : arr) std::cout << x << " ";
    std::cout << "\n";
    // ===== GPU kernel 调用 =====
    float* d_arr;
    cudaMalloc(&d_arr, 5 * sizeof(float));
    kernel_fill<<<1, 5>>>(d_arr, 5);
    cudaMemcpy(arr, d_arr, 5 * sizeof(float), cudaMemcpyDeviceToHost);
    std::cout << "After kernel_fill: ";
    for (auto x : arr) std::cout << x << " ";
    std::cout << "\n";
    cudaFree(d_arr);
    return 0;
}

# CMakeLists.txt — 现代方式（推荐）
cmake_minimum_required(VERSION 3.18) # 3.18+ 更好，3.10 也可
# CUDA 架构：可设置为 AUTO 或指定 compute capability 列表，例如 70;75;80
if(NOT DEFINED CUDA_ARCHS)
  set(CUDA_ARCHS native)  
endif()
set(CMAKE_CUDA_ARCHITECTURES ${CUDA_ARCHS})
project(MyCudaProject LANGUAGES CXX CUDA)
# C/C++ 标准
set(CMAKE_CXX_STANDARD 20)
set(CMAKE_CXX_STANDARD_REQUIRED ON)
set(CMAKE_CUDA_STANDARD 20)
set(CMAKE_CUDA_STANDARD_REQUIRED ON)
# 编译选项（示例）
add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:-Xcompiler=-fPIC>")
add_compile_options("$<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>")
# 查找源码（修改成你自己的源文件）
set(SRC
    main.cu
)
# 可选：把 CUDA separable compilation 打开（如果有 device 函数跨文件调用）
set(CMAKE_CUDA_SEPARABLE_COMPILATION ON)
# 可选：如果要生成位置无关代码（对于 shared lib 有用）
set(CMAKE_POSITION_INDEPENDENT_CODE ON)
# 创建可执行文件 / 库
add_executable(my_app ${SRC})
# 如果你的 CUDA 运行时需要链接（通常 CMake 会自动处理）
find_package(CUDAToolkit REQUIRED) # modern CMake module
target_link_libraries(my_app PRIVATE CUDA::cudart)
# 包含目录
target_include_directories(my_app PRIVATE ${CMAKE_SOURCE_DIR}/include)
# 目标属性（示例）
target_compile_features(my_app PRIVATE cxx_std_20)
# 如果你想让 nvcc 使用特定编译选项：
target_compile_options(my_app PRIVATE
  "$<$<COMPILE_LANGUAGE:CUDA>:--expt-relaxed-constexpr>"
)
# 安装规则（可选）
install(TARGETS my_app RUNTIME DESTINATION bin)

Before fill_ones (vector): 1 2 3 4 5 
After fill_ones (vector): 1 1 1 1 1 
After fill_ones (array): 1 1 1 1 1 
After kernel_fill: 1 1 1 1 1 

1⃣ 背景问题

你最开始的函数是：

template<typename Container>
__host__ __device__ constexpr
void fill_ones(Container & ct) {
    for (auto & x : ct) x = 1;
}

在调用 std::vector<int> 时：

std::vector<int> v{1,2,3,4,5};
fill_ones(v);

会出现 CUDA 编译器警告：

warning #20014-D: calling a __host__ function from a __host__ __device__ function is not allowed

原因：

• __host__ __device__ 表示函数可以在 CPU 和 GPU 调用。
• std::vector 是 host-only 类型，不能在 GPU device 函数里调用。
  数学上可以抽象成类型约束映射：
  $${\text{func}}_{\text{target}}:T\mapsto \{\begin{array}{ll}\text{host-only},& T\in H\\ \text{device-only},& T\in D\\ \text{host+device},& T\in H\cap D\end{array}$$
  这里 $H$ 是 host-only 类型集合，$D$ 是 device-compatible 类型集合。

2⃣ Host/Device 模板宏解决方案

(a) 三重模板派发宏

#define MACRO( targ_, hdc_, func_ ) \
template< targ_, HDC x = hdc_ > \
requires( x == HDC::Hst ) __host__ func_ \
template< targ_, HDC x = hdc_ > \
requires( x == HDC::Dev ) __device__ func_ \
template< targ_, HDC x = hdc_ > \
requires( x == HDC::HstDev ) __host__ __device__ func_

用法示例：

MACRO( typename Container,
       hdc<Container>,
       void fill_ones(Container & ct) { for(auto & x : ct) x = 1; } )

说明：

• 宏生成 三个版本的函数模板：
  1. Host-only：__host__
  2. Device-only：__device__
  3. Host+Device：__host__ __device__
• hdc<T> 是类型 trait，根据类型选择派发目标。

(b) Dispatcher 包装函数

#pragma nv_exec_check_disable
template<typename T>
__host__ __device__ void wrap_impl(T&& t) {
    t.func();
}
function_dispatch_macro((typename T), (hdc<T>),
                        (void wrap(T&& t)),
                        (return wrap_impl(std::forward<T>(t));))

• Dispatcher 自动根据类型实例化对应版本的 host/device 函数。
• 宏展开可避免重复写模板。
• #pragma nv_exec_check_disable 可以抑制一些 CUDA host/device 警告。
  数学表示：
  $$\text{wrap}(T)={\text{func}}_{\text{hdc}(T)}(T)$$
• CPU 调用 → 调用 host 版本
• GPU 调用 → 调用 device 版本

3⃣ 例子：兼容 Host/Device 的数组类型

template<typename T, int N>
struct hstdev_array {
    static constexpr HDC hdc = HDC::HstDev;
    __host__ __device__ T& operator[](int idx) { return data_[idx]; }
    T data_[N];
};
__global__ void kernel() {
    hstdev_array<int, 5> hd;
    fill_ones(hd); // GPU 上安全
}
int main() {
    std::array<int,5> a{1,2,3,4,5};
    fill_ones(a); // CPU 上安全
}

说明：

• hstdev_array 是 host+device 兼容数组。
• std::array 默认 HDC::Hst，只能在 CPU 调用。
• Dispatcher + 宏派发保证 类型安全，避免 host-only 调用进入 device 代码。

4⃣ CUDA Proposal（未来扩展）

1. Conditional host/device annotation：让 __host__ 和 __device__ 接受布尔参数，例如：

template<typename T>
__host__(hdc<T>==HDC::Hst)
__device__(hdc<T>==HDC::Dev)
void wrap() {
    T{}.func();
}

2. Forbid bad cross function calls：编译器保证 host-only 函数绝不在 device 上被调用，避免手动宏和模板技巧。

5⃣ 总结

1. 直接在 __host__ __device__ 函数里调用 host-only 类型（如 std::vector）会报错。
2. 使用 三重模板派发宏 + trait 可以安全生成 host/device 版本函数。
3. Dispatcher + wrapper 保证了 CPU/GPU 统一接口：
   $$\forall T,\text{wrap}(T)={\text{func}}_{\text{hdc}(T)}(T)$$
4. 将来 CUDA 扩展可能直接允许 条件 host/device 注解，无需宏。