Come posso recuperare le informazioni sull'elenco dei parametri per un kernel CUDA 4.0+?

Question

In base alla documentazione NVidia per il cuLaunchKernel function, i kernel compilati con CUDA 3.2+ contengono informazioni relative all'elenco dei parametri. C'è un modo per recuperare queste informazioni a livello di codice da un handle CUfunction? Ho bisogno di conoscere il numero di argomenti e la dimensione di ogni argomento in byte di un kernel dal suo handle CUfunction. Ho visto la documentazione NVidia di cui sopra che afferma che questa informazione esiste, ma non ho visto da nessuna parte nella documentazione CUDA che indica un modo programmatico per accedere a queste informazioni.

Per aggiungere ulteriori spiegazioni: sto lavorando con un sistema middleware. La sua libreria frontside sostituisce libcuda (la libreria API del driver) sul sistema di destinazione. La parte posteriore viene quindi eseguita come daemon su un altro host che utilizza la risorsa GPGPU e chiama la vera libcuda su quella macchina. Esistono altre soluzioni middleware che già fanno questo con cuLaunchKernel, quindi è sicuramente possibile. Inoltre, CUDA stesso usa queste informazioni per sapere come analizzare i parametri dal puntatore che passi in cuLaunchKernel.

Modifica: Originariamente avevo la versione CUDA in cui sono stati introdotti questi metadati elencati in modo errato. Era 3.2, non 4.0, secondo lo cuLaunchKernel documentation.

Answer 1

cuLaunchKernel è progettato per avviare i kernel per i quali si conosce il prototipo della funzione. Non esiste API per "reverse engineering" il prototipo di funzione.

Answer 2

Sto lavorando sullo stesso problema (non so se nel frattempo l'hai risolto). Sto usando un kernel noto per indagare su come viene usata la memoria puntata di CUfunction. Questa è la versione senza parametri:

#include<cstdio> 

extern "C" { 
    __global__ void HelloWorld(){ 
     int thid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    } 
} 

Questa è la versione un parametro e così via.

#include<cstdio> 

extern "C" { 
    __global__ void HelloWorld(int a) { 
     int thid = (blockIdx.x * blockDim.x) + threadIdx.x; 
    } 
} 

Ti suggerisco di scaricare i primi 1024 byte della memoria indicata da CUfunction e seguire i puntatori. Ad esempio all'offset 0x30 c'è un puntatore che punta a una tabella di puntatori. Ho notato che la dimensione della struttura inviata da CUfunction non cambia con il numero dei parametri di funzione, quindi la tabella che stiamo cercando deve essere cacciata seguendo i puntatori.