CUDA 内核不重叠

Question

我有一个简单的向量乘法内核，我正在为 2 个流执行它。但是当我在 NVVP 中配置文件时，内核似乎没有重叠。是不是因为每个内核执行都使用了 %100 的 GPU，如果不是，可能是什么原因？

源代码：

#include "common.h"
#include <cstdlib>
#include <stdio.h>
#include <math.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "cuda_profiler_api.h"
#include <string.h>

const int N = 1 << 20;

__global__ void kernel(int n, float *x, float *y)
{
    int i = blockIdx.x*blockDim.x + threadIdx.x;
    if (i < n) y[i] = x[i] * y[i];
}

int main()
{

    float *x, *y, *d_x, *d_y, *d_1, *d_2;
    x = (float*)malloc(N*sizeof(float));
    y = (float*)malloc(N*sizeof(float));

    cudaMalloc(&d_x, N*sizeof(float));
    cudaMalloc(&d_y, N*sizeof(float));
    cudaMalloc(&d_1, N*sizeof(float));
    cudaMalloc(&d_2, N*sizeof(float));

    for (int i = 0; i < N; i++) {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, y, N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_1, x, N*sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_2, y, N*sizeof(float), cudaMemcpyHostToDevice);

    const int num_streams = 8;

    cudaStream_t stream1;
    cudaStream_t stream2;

    cudaStreamCreateWithFlags(&stream1, cudaStreamNonBlocking);
    cudaStreamCreateWithFlags(&stream2, cudaStreamNonBlocking);

    cudaEvent_t start, stop;
    float elapsedTime;

    cudaEventCreate(&start);
    cudaEventRecord(start, 0);

    for (int i = 0; i < 300; i++) {
        kernel << <512, 512, 0, stream1 >> >(N, d_x, d_y);
        kernel << <512, 512, 0, stream2 >> >(N, d_1, d_2);
    }

    cudaStreamSynchronize(stream1);
    cudaStreamSynchronize(stream2);
    // cudaDeviceSynchronize();

    cudaEventCreate(&stop);
    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&elapsedTime, start, stop);
    printf("Elapsed time : %f ms\n", elapsedTime);

    cudaDeviceReset();
    cudaProfilerStop();
    return 0;
}

编辑：从 cmets 我了解到每个内核都在充分利用 GPU，那么实现 262144 大小的向量乘法（用于多个流）的最佳方法是什么？

我的设备信息：

CUDA Device Query...
There are 1 CUDA devices.

CUDA Device #0
Major revision number:         5
Minor revision number:         0
Name:                          GeForce GTX 850M
Total global memory:           0
Total shared memory per block: 49152
Total registers per block:     65536
Warp size:                     32
Maximum memory pitch:          2147483647
Maximum threads per block:     1024
Maximum dimension 0 of block:  1024
Maximum dimension 1 of block:  1024
Maximum dimension 2 of block:  64
Maximum dimension 0 of grid:   2147483647
Maximum dimension 1 of grid:   65535
Maximum dimension 2 of grid:   65535
Clock rate:                    901500
Total constant memory:         65536
Texture alignment:             512
Concurrent copy and execution: Yes
Number of multiprocessors:     5
Kernel execution timeout:      Yes

Answer 1

您的内核不重叠的原因是因为您的 gpu 被@Robert Crovella 提到的执行线程“填充”。检查CUDA Programming Guide 中的计算能力一章，对于您的 CC (5.0)，每个 SM 有 2048 个线程的限制。你有 5 个 SM，所以这就成功了 最多可以在您的设备上同时运行 10240 个线程。您正在调用 512x512=262144 个线程，只需要一个内核调用，这几乎没有为其他内核调用留下任何空间。

您需要启动足够小的内核，以便 2 个内核可以在您的设备上同时运行。

我不是流方面的专家，但据我了解，如果您想使用流运行程序，则需要将其拆分成块，并且必须计算适当的偏移机制才能您的流能够访问其正确的数据。在您当前的代码中，您启动的每个流对完全相同的数据执行完全相同的计算。您必须在流之间拆分数据。

除此之外，如果您想获得最大性能，您需要将内核执行与异步数据传输重叠。最简单的方法是为您的每个流分配一个如下所示的方案，如呈现的here

for (int i = 0; i < nStreams; ++i) {
     int offset = i * streamSize;
     cudaMemcpyAsync(&d_a[offset], &a[offset], streamBytes,        cudaMemcpyHostToDevice, stream[i]);
     kernel<<<streamSize/blockSize, blockSize, 0, stream[i]>>>(d_a, offset);
     cudaMemcpyAsync(&a[offset], &d_a[offset], streamBytes, cudaMemcpyDeviceToHost, stream[i]);
}

此配置只是告诉每个流执行 memcpy，然后对一些数据执行内核，然后将数据复制回来。在异步调用之后，流将同时工作以完成它们的任务。

PS：我也建议修改你的内核。使用一个线程只计算一个乘法是一种过度杀伤。我会使用线程来处理更多数据。