Les noyaux CUDA ne se chevauchent pas

Question

J'ai un simple noyau de multiplication vectorielle, que j'exécute pour 2 flux. Mais quand je profile dans NVVP, les noyaux ne semblent pas se chevaucher. Est-ce parce que chaque exécution de noyau utilise% 100 de GPU, sinon quelle peut être la cause?

code source:

#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: D'après les commentaires que je comprends chaque noyau est en utilisant pleinement GPU, alors quelle est la meilleure approche pour atteindre 262144 taille multiplication du vecteur (pour plusieurs cours d'eau) ?

Mes informations sur le périphérique:

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

La raison pour laquelle vos noyaux ne se chevauchent pas parce que votre gpu est « remplie » avec des fils d'exécution comme Crovella mentionne @ Robert. En cochant le chapitre Compute Capabilities du CUDA Programming Guide, il y a une limite de 2048 threads par SM pour votre CC (5.0). Vous avez 5 SM donc cela fait un maximum de 10240 discussions qui peuvent s'exécuter simultanément sur votre appareil. Vous appelez 512x512 = 262144 threads, avec juste un seul appel de noyau, et cela ne laisse pratiquement aucun espace pour l'autre appel du noyau.

Vous devez lancer des noyaux suffisamment petits pour que 2 puissent fonctionner simultanément sur votre périphérique. Je ne suis pas un expert sur les flux, mais d'après ce que j'ai compris, si vous voulez exécuter votre programme en utilisant des flux, vous devez le diviser en morceaux et vous devez calculer un mécanisme de décalage approprié pour que vos flux puissent accéder à leurs propres données. Sur votre code actuel, chaque flux que vous lancez effectue exactement le même calcul sur exactement les mêmes données. Vous devez diviser les données entre les flux. A part cela, si vous voulez obtenir la performance maximale, vous devez chevaucher l'exécution du noyau avec les transferts de données asynchrones. Le moyen le plus simple est d'attribuer un système comme suit pour chacun de vos flux comme présenté 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]); 
} 

Cette configuration indique simplement chaque flux de faire un memcpy puis d'exécuter le noyau sur certaines données puis de copier les données de retour. Après les appels asynchrones, les flux fonctionneront simultanément en accomplissant leurs tâches. PS: Je recommanderais également de réviser votre noyau aussi bien. Utiliser un thread pour calculer une seule multiplication est une surcharge. Je voudrais utiliser le fil pour traiter plus de données.