CuPY：没有看到内核并发

Question

我目前正在使用 CuPY 的 RawKernels 使用异步流并行化巨大的矩阵计算。

似乎每个 RawKernel 调用都在等待前一个内核完成，尽管我指定该流是非阻塞的。

有人知道我做错了什么吗？

这是一个创建 32 个流的简单示例。每个流应将 3D 输入数组的单个切片复制到 3D 输出数组。

import cupy

kernel = cupy.RawKernel(
    '''
    extern "C"
    __global__ void simple_copy(float* iArr, float* oArr, int rows, int cols, int slice){
        unsigned int col = blockDim.x*blockIdx.x + threadIdx.x;
        unsigned int row = blockDim.y*blockIdx.y + threadIdx.y;
    
        if(row < rows && col < cols){
//this for loop is just additional work to see kernel launches in visual profiler more easily
            for(int i=0; i<1000; i++){
                oArr[rows*cols*slice + row*cols + col] = iArr[rows*cols*slice + row*cols + col];
            }
        }        
    } 
    '''
    , 'simple_copy')


device = cupy.cuda.Device()
# [x, y, z]
iArr1 = cupy.ones((32*32, 32*32, 32), dtype=cupy.float32)
oArr1 = cupy.zeros((32*32, 32*32, 32), dtype=cupy.float32)

n = 32
map_streams = []
for i in range(n):
    map_streams.append(cupy.cuda.stream.Stream(non_blocking=True))

# I want to run kernel on individual z-axis slice asynchronous
for i, stream in enumerate(map_streams):
    with stream:
        kernel((32, 32), (32, 32), (iArr1, oArr1, 32*32, 32*32, i))
device.synchronize()

Answer 1

It seems like each RawKernel call is waiting for prevous kernel to finish eventhough I specify that stream is non-blocking.... .... Does anyone have an idea on what I'm doing wrong?

你并没有做错任何事，除了期待一些不可能发生的事情。

只有当有足够的资源可同时用于运行多个内核时，才有可能并发内核执行。所有当前支持的 GPU 每个多处理器最多有 2048 个活动线程，每个块有 1024 个线程。这意味着每个多处理器最多可以运行两个块。根据您的 GPU 的大小，这意味着少于大约 60 个块可以运行同时达到绝对最大值。鉴于一次内核启动会使您的 GPU 容量“饱和”多次，因此第二次实际拥有运行资源的可能性非常小。这就是为什么您看不到内核启动之间没有重叠或并发。

CuPY：没有看到内核并发

CuPY: Not seeing kernel concurrency

python

cuda

cupy