Ядра 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; 
} 

EDIT: Из комментариев я понимаю каждое ядро ​​полностью использует 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

Причина, почему ваши ядра не перекрывают друг друга, потому что ваша видеокарта «заполнена» с исполнением нитями, как упоминает @Robert Crovella. При проверке главы «Возможности вычислений» из CUDA Programming Guide существует ограничение на 2048 потоков на SM для вашего CC (5.0). У вас есть 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]); 
} 

Эта конфигурация просто говорит каждый поток, чтобы сделать тетсру затем выполнить ядро ​​на некоторых данных, а затем скопировать данные обратно. После асинхронных вызовов потоки будут работать одновременно, выполняя свои задачи.

PS: Я также рекомендовал бы пересмотреть ваше ядро. Использование одного потока для вычисления только одного умножения является излишним. Я бы использовал поток для обработки еще нескольких данных.