Tensorflow: was bedeutet Index in CUDA_1D_KERNEL_LOOP (index, nthreads) op Benutzer

Question

Ich habe in mehreren Standard-Ops von Tensorflow-Schichten (wie https://github.com/tensorflow/tensorflow/blob/master/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc), den Code CUDA_1D_KERNEL_LOOP (Index, nthreads) als Teil der Vorwärts-und Rückwärtsdurchläufe gesehen ...

Ich denke, dass der "Index" hier irgendwie auf die Koordinaten der unteren Feature Map bezogen ist, aber ich bin mir seiner genauen Bedeutung nicht sicher ... Wer könnte helfen?

Answer 1

CUDA_1D_KERNEL_LOOP(i, n) ist ein Präprozessor-Makro, das in tensorflow/core/util/cuda_kernel_helper.h definiert ist. Es stellt eine allgemeine Steuerflussanweisung bereit, die in vielen Cuda-Kernen innerhalb der Tensorflow-Codebasis verwendet wird.

Die Anweisung wird normalerweise zum Durchlaufen von Elementen eines Arrays innerhalb eines Kernels verwendet. Das Argument i ist der Name der Steuervariablen und das Argument n ist die Abbruchbedingung für die Steueranweisung. Cuda-Kernel werden in parallelen Threads gestartet. Jeder Thread arbeitet typischerweise mit einer Teilmenge von Array-Elementen. Das Makro bietet einen gewissen Komfort für den Zugriff auf die gewünschten Array-Elemente.

In dem Beispiel, das Sie verknüpfen, CUDA_1D_KERNEL_LOOP(index, nthreads) als interpretiert:

for (int index = blockIdx.x * blockDim.x + threadIdx.x; index < nthreads; index += blockDim.x * gridDim.x) 

Daher index deklariert und initialisiert innerhalb CUDA_1D_KERNEL_LOOP vor dem nachfolgenden Codeblock eingeben. Die genaue Bedeutung von index hängt davon ab, wie es innerhalb des Codeblocks verwendet wird.

Answer 2

Eine Sache, die mich verwirrte, wenn ich dieses Makro zum ersten Mal lese, ist: "Warum ist es eine Schleife, ist das nicht innerhalb des Kernels, der bereits parallelisiert wurde?" Die Antwort ist, die Schleife behandelt den Fall, wenn Sie mehr "Threads" haben, als Ihre GPU tatsächlich unterstützt. Wenn Sie beispielsweise eine parallelisierte Vektoraddition durchführen, haben Sie entschieden, dass Sie für Ihre GPU 512 Threads pro Block verwenden und maximal 4096 Blöcke planen (dies sind die Standardparameter in Caffe2). Dies bedeutet, dass Sie nur maximal 2097152 Threads einplanen können. Angenommen, Ihr Vektor hat tatsächlich 4M Elemente; Jetzt können Sie keinen Thread pro Element zuweisen. Also muss jeder Thread dafür verantwortlich sein, mehr als ein Element im Vektor zu summieren: dafür ist diese Schleife gedacht!

Hier ist ein kleineres Beispiel, das genau beschreibt, wie die Arbeit geplant wird. Angenommen, dass blockDim.x == 2, gridDim.x == 2 und nthreads == 7. Wenn wir dann einen GPU-Thread als (blockIdx.x, threadIdx.x) identifizieren, weisen wir sie zu, um die folgenden Arbeiten an dem Vektor auszuführen: [(0,0), (0,1), (1,0), (1,1), (0,0), (0,1), (1,0)]. Insbesondere können wir sehen, dass gemäß der Rastergröße nur vier GPU-Threads verfügbar sind; so für blockIdx.x == 0threadIdx.x == 0, index wird Verarbeitung der Vektorelemente bei BEIDE 0 und 4.