NVCC - Verschiedene Blockgrößen, je nach Bogen zur Kompilierungszeit

Question

Ich habe einen Kernel, der die höchste Leistung für die höchste Leistung zeigt anders Blockieren Sie die Größen beim Ausführen von Kepler und Fermi -Hardware. Ich möchte, dass ich zum Zeitpunkt der Kompilierung das aktuelle Architekturziel überprüfen und a definieren kann THREADS_PER_BLOCK Makro zu i) Starten Sie den Kernel mit; ii) die Anzahl der erforderlichen Blöcke bestimmen; iii) Stellen Sie die gemeinsame Speichergröße statisch im Kernel ein.

Das folgende zeigt, was ich versuche zu tun. Angenommen, ich ziele auf GK104 -Hardware und somit verwendet nvcc -arch=sm_30. Dies wird immer noch dazu führen THREADS_PER_BLOCK = 256 seit __CUDA_ARCH__ ist nicht für die Host -Code -Kompilierung definiert. (Ich verstehe, aus z. Diese Antwort, warum es nicht so funktionieren kann.)

#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;
}

Ich könnte die Eigenschaften der Geräte zur Laufzeit überprüfen und dynamisch gemeinsam genutzten Speicher verwenden, möchte jedoch wissen, ob dies bei der Kompilierungszeit hart codiert werden kann, ohne dass z. B. manuell einen -dfermi oder -dkepler und Einstellung hinzufügen muss THREADS_PER_BLOCK darauf bezogen. NB: Alle Benutzer dieses Codes werden ihn selbst zusammenstellen, fast Sicherlich für eine Architektur ist dies keine unangemessene Option. Es scheint nur überflüssig, wenn man an der Übergabe des -arch= Flagge.

Answer 1

nvcc Der Compiler erkennt nicht lokal verfügbare GPUs, sondern zielt standardmäßig immer auf SM 1.0 ab. Andernfalls könnte es ein ziemlich verwirrendes Verhalten beim Aufbau verschiedener Systeme einführen.

Um für das verfügbare Gerät zu kompilieren, müssen Sie den Benutzer entweder auffordern, die SM -Version anzugeben oder einen Erkennungscode während der Build -Zeit auszuführen. Ich bin überzeugt, dass es einfacher ist, Hardware -Erkennungscode in die Laufzeit zu bringen und dann Ihren Kernel -Start nach Wunsch zu konfigurieren.