numba.cuda.jit 编译的函数不强制要求 device=True，但用途不同：不加时生成 host-side kernel，可从 CPU 启动；加 device=True 才生成仅 GPU 内调用的 device function，无返回值且不可调用 host 函数。
numba.cuda.jit 编译的函数必须用 device=True 吗？
不是必须，但绝大多数情况要设 device=True。不加这个参数，默认生成的是 host-side 可调用的 kernel（即从 CPU 启动执行），它不能被其他 GPU 函数直接调用；加了 device=True 才生成 device function（类似 CUDA 的 __device__ 函数），只能在 GPU 上被其他 GPU 函数调用，不能从 CPU 直接 launch。
常见错误现象：TypingError: Failed in nopython mode pipeline 或提示 “cannot be called from cpu context”，往往就是混用了 host kernel 和 device function。
- 想从 Python 脚本里直接
my_kernel[grids, blocks](...)—— 用@cuda.jit（不加device=True） - 想在另一个
@cuda.jit函数内部调用辅助计算逻辑 —— 辅助函数必须加@cuda.jit(device=True) - device function 不能有返回值（返回类型必须是
void），也不能调用 host 函数（如print、len）
如何正确传入数组并保证内存位于 GPU 上？
Numba CUDA 不会自动把 NumPy 数组搬上 GPU；你得显式用 cuda.to_device() 搬运，否则运行时会报 CudaAPIError: Call to cuMemcpyHtoD results in CUDA_ERROR_INVALID_VALUE。
使用场景：输入是 np.ndarray，你想在 GPU kernel 里读写它。
立即学习“Python免费学习笔记（深入）”；
- 先用
gpu_arr = cuda.to_device(host_arr)上传；host_arr 必须是 C-contiguous（可用np.ascontiguousarray()强制） - kernel 参数接收的是
cuda.devicearray.DeviceNDArray，不是普通 NumPy 数组 - 结果要拿回 CPU，必须显式调用
gpu_arr.copy_to_host()；不调就不会同步，可能拿到旧数据或垃圾值 - 避免反复上传下载：对多次迭代的计算，上传一次、kernel 多次调用、最后下载一次
grid 和 block 的尺寸怎么设才不越界？
GPU kernel 里常用 cuda.grid(1) 或 cuda.grid(2) 获取全局线程索引，但它不会自动保护越界访问 —— 如果 grid 尺寸太大，线程数超过数组长度，就会读写非法内存，表现可能是静默错误、NaN、或直接 cudaErrorLaunchFailure。
典型写法是手动检查：
@cuda.jit
def my_kernel(arr):
idx = cuda.grid(1)
if idx < arr.size: # ← 这行不能少
arr[idx] = arr[idx] * 2
- 一维数组常用
blocks = (arr.size + TPB - 1) // TPB，其中TPB = 256（threads per block）是较稳妥的默认值 - 二维数组用
cuda.grid(2)时，同样要分别判断i 和 <code>j - 别盲目设大 grid：NVIDIA GPU 通常单 kernel 最多支持约 65535 × 65535 个 block，超出会启动失败
为什么我的 kernel 运行没报错但结果不对？
最常被忽略的是同步问题：CPU 发起 kernel 启动后立即往下走，不会等 GPU 执行完。如果紧接着就调用 copy_to_host()，很可能拿到未更新的数据。
- 必须加
cuda.synchronize()（或等价地，在copy_to_host()前调用）来阻塞 CPU，直到所有前序 kernel 完成 - 也可以用
stream = cuda.stream()+stream.synchronize()做细粒度控制，但新手建议先用全局同步 - 另一个隐蔽坑：CUDA kernel 中不能用
math.sqrt，得用cuda.math.sqrt（Numba 提供的设备版数学函数）；用错会导致静默降级为 0 或 NaN - 调试时可临时加
cuda.atomic.add(arr, 0, 1)统计实际执行线程数，验证是否所有线程都跑到了
device function 的调用栈深度、共享内存用量、寄存器压力这些，只有在真正遇到 launch failure 或性能瓶颈时才需要深挖 —— 初期先把数据搬运、索引边界、同步这三件事做对，就已经覆盖了 90% 的硬性错误。
