CuPY：没有看到内核并发答案

【问题标题】：CuPY: Not seeing kernel concurrencyCuPY：没有看到内核并发
【发布时间】：2020-07-18 22:22:35
【问题描述】：

我目前正在使用 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()

【问题讨论】：

你的每个内核都会填满 GPU，没有空间让其他内核执行。因此它们序列化。没有理由假设内核自动/始终同时运行。内核并发需要非常具体的条件。你还没有满足这些条件。您正在启动 1024 个块，每个 1024 个线程。 Tesla V100 一次最多只能处理 160 个这些块。没有理由假设第二个内核的某些块会与第一个内核混合，即使这样做也没有任何好处。 GPU 没有无限容量。
@RobertCrovella 感谢您的见解。在新的 CUDA 编程中，还有很多文档需要阅读。我用较小的数组测试了相同的程序并得到了正确的“并发”行为。您可以将此评论作为答案发布，以便我接受吗？
@RobertCrovella 您能分享一下您是如何从分析器创建图形输出的吗？ Cupy 关于此事的文件仍在制作中。谢谢。
@bernard 我认为您标记了错误的人。我没有在这里创建任何分析器输出。

标签： python cuda cupy

【解决方案1】：

似乎每个 RawKernel 调用都在等待前一个内核完成，即使我指定流是非阻塞的...... .... 有人知道我做错了什么吗？

除了期望不可能发生的事情之外，您没有做错任何事情。

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

【讨论】：