CUDA 12.1 - Ora supporta large parameters kernel

Post Views: 841

La release CUDA 12.1 di Nvidia aggiunge una caratteristica a lungo richiesta dagli sviluppatori: il passaggio di parametri più grandi di 4kB.

I parametri delle funzioni del kernel CUDA vengono passati al dispositivo tramite la memoria costante e sono limitati a 4.096 byte. CUDA 12.1 aumenta questo limite di parametro da 4.096 byte a 32.764 byte su tutte le architetture di dispositivi, inclusa NVIDIA Volta e versioni successive.

In precedenza, il passaggio di argomenti del kernel superiori a 4.096 byte richiedeva di aggirare il limite dei parametri del kernel copiando gli argomenti in eccesso nella memoria costante con cudaMemcpyToSymbol o cudaMemcpyToSymbolAsync, come mostrato nello snippet di seguito.

#define TOTAL_PARAMS        (8000) // ints
#define KERNEL_PARAM_LIMIT  (1024) // ints
#define CONST_COPIED_PARAMS (TOTAL_PARAMS - KERNEL_PARAM_LIMIT)

__constant__ int excess_params[CONST_COPIED_PARAMS];

typedef struct {
    int param[KERNEL_PARAM_LIMIT];
} param_t;

__global__ void kernelDefault(__grid_constant__ const param_t p,...) {
    // access <= 4,096 parameters from p
    // access excess parameters from __constant__ memory
}

int main() {
    param_t p;    
    int *copied_params = (int*)malloc(CONST_COPIED_PARAMS * sizeof(int));
    cudaMemcpyToSymbol(excess_params,
                       copied_params, 
                       CONST_COPIED_PARAMS * sizeof(int), 
                       0,
                       cudaMemcpyHostToDevice);
    kernelDefault<<<GRIDDIM,BLOCKDIM>>>(p,...);
    cudaDeviceSynchronize();
}

#define TOTAL_PARAMS (8000) // ints

#define KERNEL_PARAM_LIMIT (1024) // ints

#define CONST_COPIED_PARAMS (TOTAL_PARAMS - KERNEL_PARAM_LIMIT)

__constant__ int excess_params[CONST_COPIED_PARAMS];

typedef struct {

int param[KERNEL_PARAM_LIMIT];

} param_t;

__global__ void kernelDefault(__grid_constant__ const param_t p,...) {

// access <= 4,096 parameters from p

// access excess parameters from __constant__ memory

}

int main() {

param_t p;

int *copied_params = (int*)malloc(CONST_COPIED_PARAMS * sizeof(int));

cudaMemcpyToSymbol(excess_params,

copied_params,

CONST_COPIED_PARAMS * sizeof(int),

cudaMemcpyHostToDevice);

kernelDefault<<<GRIDDIM,BLOCKDIM>>>(p,...);

cudaDeviceSynchronize();

}

Questo approccio limita l’usabilità perché è necessario gestire in modo esplicito tanto l’allocazione di memoria costante quanto la copia. L’operazione di copia aggiunge inoltre una latenza significativa, riducendo le prestazioni dei kernel vincolati alla latenza che accettano parametri superiori a 4.096 byte.

A partire da CUDA 12.1, è possibile passare fino a 32.764 byte come parametri del kernel su NVIDIA Volta e versioni successive, risultando nell’implementazione semplificata mostrata nel secondo frammento di seguito.

#define TOTAL_PARAMS (8000) // ints

typedef struct {
    int param[TOTAL_PARAMS];
} param_large_t;

__global__ void kernelLargeParam(__grid_constant__ const param_large_t p,...) {
    // access all parameters from p
}

int main() {
    param_large_t p_large;
    kernelLargeParam<<<GRIDDIM,BLOCKDIM>>>(p_large,...);
    cudaDeviceSynchronize();
}

#define TOTAL_PARAMS (8000) // ints

typedef struct {

int param[TOTAL_PARAMS];

} param_large_t;

__global__ void kernelLargeParam(__grid_constant__ const param_large_t p,...) {

// access all parameters from p

}

int main() {

param_large_t p_large;

kernelLargeParam<<<GRIDDIM,BLOCKDIM>>>(p_large,...);

cudaDeviceSynchronize();

}

In entrambi gli esempi precedenti, i parametri del kernel sono annotati con il qualificatore __grid_constant__ per indicare che sono di sola lettura.

Per compilare, avviare ed eseguire il debug dei kernel con parametri kernel di grandi dimensioni è necessario l’uso di CUDA Toolkit 12.1 e un driver R530 o versione successiva. CUDA genererà l’errore CUDA_ERROR_NOT_SUPPORTED se si tenta l’avvio su un driver precedente.

Il limite di parametri più elevato è disponibile su tutte le architetture, inclusa NVIDIA Volta e versioni successive. Il limite del parametro rimane a 4.096 byte sulle architetture inferiori a NVIDIA Volta.

Compatibilità dei collegamenti tra le revisioni di CUDA Toolkit

Quando si collegano oggetti device, se almeno un oggetto device contiene un kernel con il limite di parametri più alto, è necessario ricompilare tutti gli oggetti dalle origini del device, linkandoli insieme con CUDA Toolkit 12.1. In caso contrario, si verificherà un errore del linker.

Consideriamo ad esempio lo scenario in cui due oggetti device, a.o e b.o, sono linkati insieme. Se a.o o b.o contengono almeno un kernel con il limite di parametri più alto, è necessario ricompilare i rispettivi sorgenti e collegare insieme gli oggetti risultanti pena il non completamento della compilazione del modulo.

Prestazioni con parametri del kernel di grandi dimensioni

La Figura mette a confronto le prestazioni dei due frammenti di codice (forniti sopra) su un singolo sistema NVIDIA H100 misurato su 1.000 iterazioni. In questo esempio, evitando copie costanti si è ottenuto un risparmio complessivo del 28% nel tempo di esecuzione dell’applicazione. Per gli stessi frammenti, la Figura 2 mostra un miglioramento del 9% nel tempo di esecuzione del kernel, misurato con NVIDIA Nsight Systems.

Consideerazioni finali

CUDA 12.1 offre la possibilità di trasferire fino a 32.764 byte utilizzando i parametri del kernel, che possono essere sfruttati per semplificare le applicazioni e ottenere miglioramenti delle prestazioni. Per vedere l’esempio di codice completo a cui si fa riferimento in questo post, visita NVIDIA/cuda-samples su GitHub.

(Fonte: Nvidia Developer Blog)

Join our groups on Telegram…

… and don’t forget our social channels!

Luigi_Morelli

Definire ciò che si è non risulta mai semplice o intuitivo, in specie quando nella vita si cerca costantemente di migliorarsi, di crescere tanto professionalmente quanto emotivamente. Lavoro per contribuire al mutamento dei settori cardine della computer science e per offrire sintesi ragionate e consulenza ad aziende e pubblicazioni ICT, ma anche perche’ ciò che riesco a portare a termine mi dà soddisfazione, piacere. Così come mi piace suonare (sax, tastiere, chitarra), cantare, scrivere (ho pubblicato 350 articoli scientfici e 3 libri sinora, ma non ho concluso ciò che ho da dire), leggere, Adoro la matematica, la logica, la filosofia, la scienza e la tecnologia, ed inseguo quel concetto di homo novus rinascimentale, cercando di completare quelle sezioni della mia vita che ancora appaiono poco ricche.