为何无法使用同一指针启动并发CUDA内核？

你的问题并非“两个内核不能使用同一指针”，而是统一内存（Unified Memory）的一致性规则与Windows平台的延迟迁移策略导致的内存访问冲突，具体分析如下：

1. 统一内存的核心特性 #

你用cudaMallocManaged分配的是统一内存，这类内存支持主机和设备直接访问，但Windows平台默认采用延迟迁移策略：只有当主机或设备实际访问某页内存时，该页才会被迁移到对应的内存空间（主机端或设备端）。同时CUDA的一致性规则明确：主机访问统一内存时，必须确保设备端对该内存的所有操作已完成，否则会触发未定义行为（比如你遇到的“error reading location”异常）。

2. 你的代码中的冲突链 #

• 调用kernelA<<<10, 10>>>(big->small)后，设备开始执行kernelA，big->small指向的内存页会被迁移到设备端并被设备占用。
• 紧接着主机代码需要读取big->small的值（作为kernelB的参数），此时主机尝试访问仍被设备占用的内存页，触发了页迁移冲突，导致内存访问错误。
• 当你添加cudaDeviceSynchronize()后，该函数会强制等待kernelA完全执行完毕，设备释放对该内存页的占用，主机才能安全访问big->small，进而正常启动kernelB。

3. 内核共用指针的可行性 #

实际上，两个内核完全可以安全地共用同一数据指针并发执行——问题的核心不是指针共享，而是主机在未同步的情况下访问了被设备占用的统一内存。如果你的big->small是普通设备指针（用cudaMalloc分配，主机不直接访问指针指向的内存），连续启动两个内核不会有任何问题，因为主机只是传递指针地址值，不需要访问指针指向的内存空间。

替代解决方案（无需同步） #

你可以提前在主机端读取big->small的地址值并保存，避免在启动第二个内核时再次访问统一内存：

void myFunc(Big* big) {
    // 提前读取统一内存中的指针值，保存到主机变量
    Small* dev_small = big->small;
    kernelA <<<10, 10>>> (dev_small);
    // 无需同步，直接传递已保存的地址值给kernelB
    kernelB <<<10, 10>>> (dev_small);
}

内容的提问来源于stack exchange，提问作者Daniel Johansen