Ich versuche, dynamische Parallelität zu verwenden, um einen Algorithmus zu verbessern, den ich in CUDA habe. In meiner ursprünglichen CUDA-Lösung berechnet jeder Thread eine Zahl, die für jeden Block gleich ist. Was ich tun möchte, ist, zuerst einen groben Kernel (oder einen Kernel mit niedriger Auflösung) zu starten, in dem Threads den gemeinsamen Wert nur einmal berechnen (als ob jeder Thread einen Block darstellt). Anschließend erstellt jeder Thread ein kleines Raster mit 1 Block (16 x 16 Threads) und startet einen untergeordneten Kernel, der den gemeinsamen Wert übergibt. Theoretisch sollte es schneller sein, weil man viele redundante Operationen spart. Aber in der Praxis funktioniert die Lösung sehr langsam, ich weiß nicht warum.

Dies ist der Code, sehr vereinfacht, nur die Idee.

__global__ coarse_kernel( parameters ){
    int common_val = compute_common_val();
    dim3 dimblock(16, 16, 1);
    dim3 dimgrid(1, 1, 1);
    child_kernel <<< dimgrid, dimblock >>> (common_val, parameters);

}

__global__ child_kernel( int common_val, parameters ){
    // use common value
    do_computations(common_val, parameters);
}

Die Anzahl der child_kernels ist sehr hoch, einer pro Thread und es müssen ungefähr 400x400 Threads vorhanden sein. Soweit ich weiß, sollte die GPU alle diese Kernel parallel verarbeiten, oder?

Oder werden Kinderkerne irgendwie sequentiell abgearbeitet?

Meine Ergebnisse zeigen, dass die Leistung mehr als zehnmal langsamer ist als in der ursprünglichen Lösung, die ich hatte.