1

我有一个简单的向量乘法内核，我正在为 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;
}

编辑：从评论中我了解到每个内核都在充分利用 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
4

1 回答 1

4

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

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

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

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

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

于 2016-02-04T21:26:08.117 回答