Ядра CUDA не перекрываются

Question

У меня есть простое ядро ​​векторного умножения, которое я выполняю для 2 потоков. Но когда я создаю профиль в NVVP, кажется, что ядра не перекрываются. Это потому, что каждое ядро ​​использует%100 графического процессора, если нет, что может быть причиной?

Исходный код:

#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

Answer 1

Причина, по которой ваши ядра не перекрываются, заключается в том, что ваш gpu "заполнен" потоками выполнения, как упоминает @Robert Crovella. В главе "Возможности вычислений" из Руководства по программированию CUDA установлено ограничение в 2048 потоков на SM для вашей CC (5.0). У вас есть 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: Я бы также порекомендовал пересмотреть ваше ядро. Использование одного потока для вычисления только одного умножения является излишним. Я бы использовал поток, чтобы обработать еще немного данных.