Array dinamici globali CUDA (come in C) allocati alla memoria del dispositivo

Question

Quindi, sto cercando di scrivere del codice che utilizzi l'architettura CUDA di Nvidia.Ho notato che la copia da e verso il dispositivo stava davvero danneggiando le mie prestazioni complessive, quindi ora sto cercando di spostare una grande quantità di dati sul dispositivo.

Poiché questi dati vengono utilizzati in numerose funzioni, vorrei che fossero globali.Sì, posso passare i puntatori, ma in questo caso mi piacerebbe davvero sapere come lavorare con i globali.

Quindi, ho funzioni del dispositivo che desiderano accedere a un array allocato del dispositivo.

Idealmente, potrei fare qualcosa del tipo:

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

Tuttavia, non ho capito come creare un array dinamico.Ho trovato una soluzione dichiarando l'array come segue:

__device__ float global_data[REALLY_LARGE_NUMBER];

E sebbene ciò non richieda una chiamata cudaMalloc, preferirei l'approccio di allocazione dinamica.

Answer 1

Qualcosa del genere probabilmente dovrebbe funzionare.

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)


__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}


int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

Provatelo.

Answer 2

Dedica un po' di tempo concentrandoti sulla copiosa documentazione offerta da NVIDIA.

Dalla Guida alla Programmazione:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

Questo è un semplice esempio di come allocare la memoria.Ora, nel kernel, dovresti accettare un puntatore a un float in questo modo:

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

Quindi ora puoi invocarli in questo modo:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

Poiché questi dati vengono utilizzati in numerose funzioni, vorrei che fosse globale.

Ci sono alcuni buoni motivi per utilizzare i globali.Questo sicuramente non lo è.Lascerò come esercizio l'espansione di questo esempio per includere lo spostamento di "devPtr" in un ambito globale.

MODIFICARE:

Ok, il problema fondamentale è questo:i tuoi kernel possono accedere solo alla memoria del dispositivo e gli unici puntatori di ambito globale che possono utilizzare sono quelli della GPU.Quando si chiama un kernel dalla CPU, dietro le quinte ciò che accade è che i puntatori e le primitive vengono copiati nei registri della GPU e/o nella memoria condivisa prima che il kernel venga eseguito.

Quindi il più vicino che posso suggerire è questo:usa cudaMemcpyToSymbol() per raggiungere i tuoi obiettivi.Ma, sullo sfondo, considera che un approccio diverso potrebbe essere la cosa giusta.

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}


int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

Non dimenticare '--host-compilation=c++' per questo esempio.

Answer 3

Sono andato avanti e ho provato la soluzione di allocare un puntatore temporaneo e passarlo a una semplice funzione globale simile al kernel1.

La buona notizia è che funziona :)

Tuttavia, penso che confonda il compilatore poiché ora ricevo "Avviso:Impossibile dire a cosa punta il puntatore, presupponendo lo spazio di memoria globale" ogni volta che provo ad accedere ai dati globali.Fortunatamente, l’ipotesi sembra essere corretta, ma gli avvertimenti sono fastidiosi.

Ad ogni modo, per la cronaca, ho esaminato molti esempi e ho eseguito gli esercizi di Nvidia in cui lo scopo è ottenere che l'output dica "Corretto!".Tuttavia non ho guardato Tutto di loro.Se qualcuno conosce un esempio di SDK in cui viene eseguita l'allocazione dinamica della memoria del dispositivo globale, mi piacerebbe comunque saperlo.

Answer 4

Ehm, il mio problema era proprio il problema di spostare devPtr in ambito globale.

Ho un'implementazione che fa esattamente questo, con i due kernel che hanno un puntatore ai dati passati.Non voglio esplicitamente passare questi puntatori.

Ho letto la documentazione abbastanza attentamente e ho visitato i forum nvidia (e ho cercato su Google per circa un'ora), ma non ho trovato un'implementazione di un array di dispositivi dinamici globale che funzioni effettivamente (ne ho provati diversi che compilano e poi fallire in modi nuovi e interessanti).

Answer 5

controlla gli esempi inclusi con l'SDK.Molti di questi progetti di esempio sono un modo decente per imparare dall’esempio.

Answer 6

Poiché questi dati vengono utilizzati in numerose funzioni, vorrei che fossero globali.

-

Ci sono alcuni buoni motivi per utilizzare i globali.Questo sicuramente non lo è.Lascerò come un esercizio per espandere questo esempio per includere lo spostamento di "Devptr" in un ambito globale.

Cosa succede se il kernel opera su una struttura const di grandi dimensioni composta da array?L'uso della cosiddetta memoria costante non è un'opzione, perché ha dimensioni molto limitate.quindi devi metterlo nella memoria globale..?