NVCC - Diverse dimensioni del blocco a seconda dell'arco al momento della compilazione

Question

Ho un kernel che mostra le prestazioni più alte per diverso Blocca le dimensioni quando si esegue Hardware Kepler e Fermi. Vorrei, a tempo di compilazione, controllare l'obiettivo di architettura attuale e definire a THREADS_PER_BLOCK Macro a i) lanciare il kernel con; ii) determinare il numero di blocchi necessari; iii) Impostare staticamente la dimensione della memoria condivisa nel kernel.

Il seguente dimostra ciò che sto tentando di fare. Supponiamo di prendere di mira l'hardware GK104 e quindi utilizzare nvcc -arch=sm_30. Questo si tradurrà ancora THREADS_PER_BLOCK = 256 da __CUDA_ARCH__ non è definito per la compilazione del codice host. (Capisco, da EG questa risposta, perché non può funzionare in questo modo.)

#if __CUDA_ARCH__ >= 300
#define THREADS_PER_BLOCK 512
#else
#define THREADS_PER_BLOCK 256
#endif

__global__ void some_kernel(int* a, int* b) {
    __shared__ sm_data[THREADS_PER_BLOCK];
    // Do something.
}

int main(void) {
    // Initialize data.
    // Calculate blocks based on THREADS_PER_BLOCK, problem size and some max.
    some_kernel<<blocks, THREADS_PER_BLOCK>>>(d_a, d_b)
    return 0;
}

Potrei controllare le proprietà del dispositivo in fase di esecuzione e utilizzare la memoria condivisa dinamica, ma vorrei sapere se questo può essere codificato al tempo di compilazione senza che EG debba aggiungere manualmente un-dfermi o -dkepler e l'impostazione THREADS_PER_BLOCK basato su questo. NB: tutti gli utenti di questo codice lo compileranno da soli, quasi Certamente per un'architettura, quindi questa non è un'opzione irragionevole. Sembra solo superfluo alla luce del passaggio del -arch= bandiera.

Answer 1

nvcc Il compilatore non rileva GPU disponibili localmente, ma mira sempre a SM 1.0 per impostazione predefinita. Altrimenti potrebbe introdurre un comportamento abbastanza confuso quando si basa su sistemi diversi.

Per compilare per il dispositivo disponibile, è necessario chiedere all'utente di specificare la versione SM o eseguire un po 'di codice di rilevamento durante il tempo di build. Sono convinto che sia più facile mettere il codice di rilevamento hardware in runtime e quindi configurare il lancio del kernel come desiderato.