Gibt es einen Geräteseitenzeiger des Hostspeichers für den Kernel in OpenCL (wie CUDA)? - Speicher, Cuda, Opencl, Null-Kopie

In CUDA können wir durch den geräteseitigen Zeiger des Hostspeichers eine vom Kernel verwaltete Datenübertragung vom Hostspeicher zum gemeinsam genutzten Speicher des Geräts erreichen. So was:

int  *a,*b,*c;          // host pointers
int *dev_a, *dev_b, *dev_c;     // device pointers to host memory

…

cudaHostGetDevicePointer(&dev_a, a, 0); // mem. copy to device not need now, but ptrs needed instead
cudaHostGetDevicePointer(&dev_b, b, 0);
cudaHostGetDevicePointer(&dev_c ,c, 0);

…

//kernel launch
add<<<B,T>>>(dev_a,dev_b,dev_c);
// dev_a, dev_b, dev_c are passed into kernel for kernel accessing host memory directly.

Im obigen Beispiel kann Kernel-Code über den Host-Speicher zugreifen dev_a, dev_b und dev_c. Kernel kann diese Zeiger verwenden, um Daten direkt vom Host in den gemeinsam genutzten Speicher zu verschieben, ohne sie über den globalen Speicher weiterzuleiten.

Aber scheint es in OpenCL eine unmögliche Mission zu sein? (Lokaler Speicher in OpenCL ist das Gegenstück zum Shared Memory in CUDA)

Antworten:

In OpenCL finden Sie genau identische API.

Wie es auf CUDA funktioniert:

Gemäß diese Präsentation und das offizielle Dokumentation.

Das Geld-Zitat über cudaHostGetDevicePointer :

Gibt den Gerätezeiger des zugeordneten Hostspeichers zurück, der von zugewiesen wurde cudaHostAlloc oder von cudaHostRegister registriert.

CUDA cudaHostAlloc mit cudaHostGetDevicePointer funktioniert genau so CL_MEM_ALLOC_HOST_PTR mit MapBuffer arbeitet in OpenCL. Wenn es sich um eine diskrete GPU handelt, werden die Ergebnisse im Gerät zwischengespeichert, und wenn es sich um eine diskrete GPU mit gemeinsamem Speicher mit dem Host handelt, wird der Speicher direkt verwendet. Es gibt also keine eigentliche "Zero Copy" -Operation mit diskreter GPU in CUDA.

Die Funktion cudaHostGetDevicePointer nimmt keine rohen malloced-Zeiger auf, genauso wieWas ist die Einschränkung in OpenCL. Aus der Sicht der API-Benutzer sind diese beiden identischen Ansätze, sodass die Implementierung nahezu identische Optimierungen durchführen kann.

Mit der diskreten GPU erhalten Sie den Zeiger auf Punkteein Bereich, in dem die GPU direkt über DMA Daten übertragen kann. Andernfalls würde der Treiber Ihren Zeiger nehmen, die Daten in den DMA-Bereich kopieren und dann die Übertragung einleiten.

Allerdings in OpenCL2.0 Dies ist explizit möglich, abhängig von den Fähigkeiten Ihrer Geräte. Mit der feinsten Granularitätsfreigabe können Sie zufällig zugewiesene Hostzeiger verwenden und sogar Atomics mit dem Host verwenden, sodass Sie sogar den Kernel vom Host aus dynamisch steuern können, während er ausgeführt wird.

http://www.khronos.org/registry/cl/specs/opencl-2.0.pdf

Auf Seite 162 finden Sie die Spezifikationen für den gemeinsam genutzten virtuellen Speicher. Beachten Sie, dass selbst wenn Sie Kernel schreiben, diese aus der Kernel-Sicht nur globale Zeiger sind.