cudaMallocManaged分配统一内存数组，指针可被CPU和GPU直接访问，无需手动拷贝；需检查返回错误、同步确保迁移完成，并注意访存模式对性能的影响。
cudaMallocManaged 分配统一内存数组最直接
统一内存（Unified Memory）在 CUDA 中本质是让 CPU 和 GPU 共享同一块逻辑地址空间，cudaMallocManaged 是分配这种内存的唯一标准接口。它返回的指针既能在主机端（CPU）读写，也能在设备端（GPU）核函数中直接访问，无需手动 cudaMemcpy。
常见错误是误用 new 或 malloc 分配普通内存后试图传给 GPU 核函数——这会触发非法内存访问或静默数据不一致；也有人混淆 cudaMalloc（纯设备内存）和 cudaMallocManaged，后者才具备自动迁移能力。
- 分配一个含 1024 个
float的统一数组：float* d_arr; cudaError_t err = cudaMallocManaged(&d_arr, 1024 * sizeof(float)); if (err != cudaSuccess) { printf("cudaMallocManaged failed: %s\n", cudaGetErrorString(err)); } - 分配后需显式调用
cudaDeviceSynchronize()或核函数启动前确保迁移完成，否则首次 GPU 访问可能触发 page fault 延迟（尤其在多 GPU 环境下） - 数组元素可直接用
d_arr[i]在主机或设备代码中访问，但注意：CPU 写入后 GPU 读取、或反之，不保证立即可见——需靠cudaStreamSynchronize或隐式同步（如核函数 launch）来保序
统一内存数组在核函数里怎么用
分配好的统一内存指针可以直接传入核函数，不需要额外处理。但要注意访问模式对性能影响极大：随机跨页访问（尤其是小步长跳转）会频繁触发迁移，远慢于连续访存。
典型场景是把数组作为输入输出参数传给 kernel，比如做向量加法：
立即学习“C++免费学习笔记（深入）”；
__global__ void add_kernel(float* a, float* b, float* c, int n) {
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx < n) c[idx] = a[idx] + b[idx];
}
<p>// 主机端调用
add_kernel<<<blocks, threads>>>(d_a, d_b, d_c, N);
cudaDeviceSynchronize(); // 必须同步，否则主机可能提前读结果- 核函数内不能调用
cudaMallocManaged或任何 CUDA 运行时 API - 统一内存指针在 kernel 中是普通指针，不支持
__restrict__优化（编译器无法假设无别名），若确定无 alias，建议复制到 local 变量再算 - 若 kernel 中只读不写，可调用
cudaMemPrefetchAsync(d_arr, size, cudaCpuDeviceId)提前将数据拉回 CPU，避免运行时缺页中断
释放统一内存必须用 cudaFree
用 cudaMallocManaged 分配的内存，必须用 cudaFree 释放，不能用 delete、free 或 cudaFreeHost —— 否则会触发 cudaErrorInvalidValue 错误，且内存泄漏。
- 释放前建议先调用
cudaDeviceSynchronize()，确保所有异步操作（包括后台迁移）已完成，避免释放中途中断迁移导致未定义行为 - 若程序有多个 GPU，且统一内存被 prefetch 到某卡上，
cudaFree仍安全，CUDA 运行时会自动清理所有设备上的副本和页表项 - 释放后指针立即失效，继续解引用会导致段错误（SIGSEGV），调试时可用
cudaPointerGetAttributes检查指针有效性（但开销大，仅限 debug）
统一内存数组不适合高频小粒度访问
统一内存不是“万能零拷贝”，它的延迟掩盖依赖硬件 page fault 和后台迁移线程。当数组被反复、小批量、跨设备交替访问（例如每轮迭代只读几个元素，然后切回 CPU 处理），实际性能可能比显式 cudaMemcpy 更差。
- 适合场景：单次加载、大量连续计算（如整张图像滤波）、算法逻辑难以拆分 CPU/GPU 阶段的原型开发
- 不适合场景：实时音视频 pipeline 中每帧只改几个像素、树形结构遍历、稀疏矩阵的非规则索引访问
- 若怀疑性能瓶颈在迁移，可用
nvidia-smi --query-compute-apps=pid,used_memory --format=csv观察内存驻留设备变化，或用nsys profile查看 page fault 次数
真正关键的是理解 unified memory 不是“自动优化”，而是“自动容错”——它保功能正确性，不保性能最优。分配数组只是第一步，访问模式和同步点设计才是决定成败的地方。
