将设备内存访问与主机线程同步

Question

CUDA内核是否可以在没有任何主机端调用(例如cudaDeviceSynchronize)的情况下同步对设备映射内存的写入？当我运行以下程序时，内核似乎不会在终止之前等待对设备映射内存的写入完成，因为在内核启动后立即检查页面锁定的主机内存不会显示内存的任何修改(除非插入延迟或取消注释对cudaDeviceSynchronize的调用)：

#include <stdio.h>
#include <cuda.h>
__global__ void func(int *a, int N) {
int idx = threadIdx.x;
if (idx < N) {
a[idx] *= -1;
__threadfence_system();
}
}
int main(void) {
int *a, *a_gpu;
const int N = 8;
size_t size = N*sizeof(int);
cudaSetDeviceFlags(cudaDeviceMapHost);
cudaHostAlloc((void **) &a, size, cudaHostAllocMapped);
cudaHostGetDevicePointer((void **) &a_gpu, (void *) a, 0);
for (int i = 0; i < N; i++) {
a[i] = i;
}
for (int i = 0; i < N; i++) {
printf("%i ", a[i]);
}
printf("n");
func<<<1, N>>>(a_gpu, N);
// cudaDeviceSynchronize();
for (int i = 0; i < N; i++) {
printf("%i ", a[i]);
}
printf("n");
cudaFreeHost(a);
}

我正在Linux上用CUDA 4.2.9为sm_20编译上述内容，并在费米GPU上运行(S2050)。

Answer 1

在发生任何内核活动之前，内核启动将立即返回到主机代码。内核执行以这种方式与主机执行异步，不会阻止主机执行。因此，您必须等待一段时间，或者使用屏障(如cudaDeviceSynchronize())来查看内核的结果，这并不奇怪。

如上所述：

为了便于主机和设备之间的并发执行，一些函数调用是异步的：控制返回到主机线程，然后设备完成请求的任务。这些是：

• 内核启动
• 内存在两个地址之间复制到同一设备内存
• 64KB或更小的内存块从主机到设备的内存拷贝
• 由以Async作为后缀的函数执行的内存复制
• 内存设置函数调用

这当然都是有意的，这样你就可以同时使用GPU和CPU。如果你不想要这种行为，你已经发现的一个简单的解决方案就是插入一个屏障。如果内核正在生成将立即复制回主机的数据，则不需要单独的屏障。内核之后的cudaMemcpy调用将等到内核完成后才开始复制操作。

我想回答你的问题，你希望内核的启动是同步的，而不需要使用屏障(你为什么要这样做？添加cudaDeviceSynchronize()调用是个问题吗？)可以这样做：

"程序员可以全局禁用所有的异步内核启动通过设置CUDA_LAUNCH_BLOCKING环境变量设置为1。此功能是仅用于调试目的，不应用作使生产软件可靠运行。">

如果您想要这种synchronous行为，最好只使用屏障(或依赖于另一个后续的cuda调用，如cudaMemcpy)。如果您使用上面的方法并依赖它，那么当其他人尝试在没有环境变量集的情况下运行它时，您的代码就会中断。所以这真的不是一个好主意。