Python(Numba)中CUDA函数进度汇报异常：始终显示0的问题

核心问题 #

1. 同流操作串行化：你把核函数和进度拷贝操作放在同一个CUDA流中，CUDA流内的操作是严格顺序执行的，拷贝操作会被排在核函数之后，导致核函数没跑完就看不到进度更新。
2. 多线程竞争写进度：256个线程同时执行大循环并修改progress[0]，会引发数据竞争，进度值可能被覆盖或更新异常。
3. 无终止条件的监控线程：监控线程会无限循环，即使核函数执行完毕也不会停止。

修复后的代码 #

import time
import numpy as np
from numba import cuda
import threading

@cuda.jit
def gpu_function(progress):
    # 仅让线程0负责更新进度，避免多线程竞争
    if cuda.threadIdx.x == 0 and cuda.blockIdx.x == 0:
        progress[0] = 17
        total_iter = 1000000000
        for i in range(total_iter):
            if i % 100000 == 0:
                progress[0] = 17 + i // 1000
                # 强制刷新全局内存，确保主机端能读到最新值
                cuda.syncthreads()
    # 其他线程可以执行实际的并行计算任务（这里示例留空）

def monitor_progress(progress_gpu, progress_host, copy_stream, stop_event):
    while not stop_event.is_set():
        # 使用独立的拷贝流，与核函数流并行执行
        progress_gpu.copy_to_host(progress_host, stream=copy_stream)
        # 等待拷贝完成再读取值
        copy_stream.synchronize()
        print(f"Progress: {progress_host[0]}")
        time.sleep(1)

progress_host = np.zeros(1, dtype=np.int32)

with cuda.pinned(progress_host):
    # 创建两个独立的流：一个跑核函数，一个跑进度拷贝
    kernel_stream = cuda.stream()
    copy_stream = cuda.stream()
    progress_gpu = cuda.to_device(progress_host, stream=kernel_stream)

    # 启动核函数（仅用一个线程块的线程0更新进度，其他线程可做计算）
    gpu_function[1, 256, kernel_stream](progress_gpu)

    # 创建停止事件，用于终止监控线程
    stop_event = threading.Event()
    monitor_thread = threading.Thread(target=monitor_progress, args=(progress_gpu, progress_host, copy_stream, stop_event))
    monitor_thread.start()

    # 等待核函数执行完毕
    kernel_stream.synchronize()
    # 通知监控线程停止
    stop_event.set()
    monitor_thread.join()
    print("GPU函数执行完成")

关键修改点 #

• 分离CUDA流：用kernel_stream执行核函数，copy_stream执行进度拷贝，两个流的操作可以并行，主机端能实时读到更新的进度。
• 单线程更新进度：仅让线程(0,0)负责进度更新，避免多线程同时写全局内存的竞争问题。
• 添加终止事件：用threading.Event控制监控线程的退出，避免核函数执行完后线程仍无限循环。
• 强制内存同步：在更新进度后调用cuda.syncthreads()，确保进度值被刷新到全局内存，主机端能读取到最新数据。

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