关于cuda.to_device异步性及流使用的技术问询

核心疑问 #

• cuda.to_device 是否为异步操作？
• cuda.to_device 是否与核函数启动使用同一个CUDA流？

根据CUDA memcpy相关文档，这类操作相对于主机是同步的，但测试结果出现矛盾点，以下是完整测试过程与疑问：

基础测试代码与结果 #

from numba import cuda
import numpy as np

A = np.ones((10000, 10000))

%timeit cuda.to_device(A)

188 ms ± 5.3 ms per loop (mean ± std. dev. of 7 runs, 10 loops each)

%timeit cuda.synchronize()

14.5 µs ± 7.03 µs per loop (mean ± std. dev. of 7 runs, 10000 loops each)

%%timeit -n1 cuda.to_device(A)
cuda.synchronize()

82.6 µs ± 11.2 µs per loop (mean ± std. dev. of 7 runs, 1 loop each)

疑问：如果cuda.to_device是同步的，为何包含同步操作的总耗时仅略高于单独同步的耗时？难道“相对于主机同步”不等于GPU已完成全部操作？

指定显式流的测试 #

stream = cuda.stream()
%timeit cuda.to_device(A, stream=stream)

188 ms ± 4.06 ms per loop (mean ± std. dev. of 7 runs, 10 loops each)

%%timeit -n1 cuda.to_device(A, stream=stream)
cuda.synchronize()

82.9 µs ± 6.66 µs per loop (mean ± std. dev. of 7 runs, 1 loop each)

疑问：按“将传输操作入队到流中”的预期，仅调用cuda.to_device应该立即返回，耗时应该更短，但实际测试并非如此。

%%timeit -n1 cuda.synchronize()
cuda.to_device(A, stream=stream)

186 ms ± 30 ms per loop (mean ± std. dev. of 7 runs, 1 loop each)

核函数流冲突测试 #

import time

start = time.time()

for i in range(1000):
  matrix_multiplication[blocks_per_grid, threads_per_block](A_gpu, B_gpu, C_gpu)

print(f'launched kernels: {time.time() - start}')
cuda.to_device(A) # 由于等待队列空闲，耗时较长
print(f'transferred: {time.time() - start}')
cuda.synchronize() 
print(f'synchronized: {time.time() - start}') # 时间几乎相同，说明cuda.to_device等待其他核函数完成

launched kernels: 0.16392898559570312
transferred: 29.859858512878418
synchronized: 29.860819101333618

推测：cuda.to_device似乎在等待其他核函数完成，说明它和核函数使用同一个流。

暂存缓冲区的疑问 #

已知CPU与GPU之间存在暂存缓冲区，to_device将数据发送到暂存缓冲区后即返回，额外延迟来自暂存缓冲区向GPU发送数据的过程。但无法理解：暂存区→GPU的耗时（82.6 µs - 14.5 µs）为何比CPU→暂存区的耗时（188 ms）快这么多？

解答 #

1. cuda.to_device的异步性本质 #

cuda.to_device并非完全同步或异步，其行为取决于主机内存类型：

• 对于可分页主机内存（如普通numpy数组）：cuda.to_device会先同步将数据拷贝到CUDA驱动管理的固定内存暂存区（这一步是主机同步的，对应你看到的188ms耗时），之后再将暂存区的数据异步拷贝到GPU设备内存——这一步会被入队到CUDA流中，主机线程无需等待即可继续执行。
• 若使用固定内存（通过cuda.pinned_array创建）：cuda.to_device会直接异步将数据从主机固定内存拷贝到GPU，此时主机线程立即返回，耗时会大幅降低。

2. CUDA流的使用逻辑 #

• 默认情况下，cuda.to_device和核函数启动都使用默认CUDA流（NULL流）。默认流的特殊之处在于：它会同步等待所有其他流的操作完成后才执行自身操作，且流内所有操作顺序执行。这就解释了大量核函数启动后调用cuda.to_device会等待核函数完成的现象——它们都在默认流里，必须按顺序执行。
• 指定显式流时，cuda.to_device的暂存区→GPU拷贝操作会入队到该显式流，但CPU→暂存区的同步拷贝仍会阻塞主机线程，所以单独调用cuda.to_device的耗时还是188ms左右；而先调用cuda.to_device再同步的测试中，CPU→暂存区已完成，同步仅需等待暂存区→GPU的异步操作，因此耗时仅略高于单独synchronize的耗时。

3. 暂存缓冲区速度差异的原因 #

• CPU→暂存区是主机普通内存到固定内存的拷贝，受限于CPU内存带宽和PCIe主机到设备方向的传输效率；而暂存区→GPU是固定内存到设备内存的直接DMA传输，由GPU硬件直接处理，无需CPU介入，且固定内存的访问延迟更低、PCIe设备方向传输效率更高，因此这一步耗时远短于CPU到暂存区的拷贝。

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