CUDA steter Erinnerung Best Practices

Präsentiere ich hier einige code

__constant__ int array[1024];

__global__ void kernel1(int *d_dst) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = array[tId];
}

__global__ void kernel2(int *d_dst, int *d_src) {
   int tId = threadIdx.x + blockIdx.x * blockDim.x;
   d_dst[tId] = d_src[tId];
}

int main(int argc, char **argv) {
   int *d_array;
   int *d_src;
   cudaMalloc((void**)&d_array, sizeof(int) * 1024);
   cudaMalloc((void**)&d_src, sizeof(int) * 1024);

   int *test = new int[1024];
   memset(test, 0, sizeof(int) * 1024);

   for (int i = 0; i < 1024; i++) {
     test[i] = 100;
   }

   cudaMemcpyToSymbol(array, test, sizeof(int) * 1024);
   kernel1<<< 1, 1024 >>>(d_array);

   cudaMemcpy(d_src, test, sizeof(int) * 1024, cudaMemcpyHostToDevice);
   kernel2<<<1, 32 >>>(d_array, d_src),

   free(test);
   cudaFree(d_array);
   cudaFree(d_src);

   return 0;
}

Das zeigt einfach nur konstant Speicher und global memory die Nutzung zu. Über die Umsetzung der "kernel2" führt zu 4-mal schneller (in Bezug auf Zeit) als "kernel1"

Verstehe ich aus dem Cuda C programming guide, dass dies, weil die Zugriffe auf constant memory sind immer serialisiert. Das bringt mich auf die Idee, dass steter Erinnerung werden am besten verwertet, wenn ein warp auf eine einzige Konstante Werte wie integer, float, double usw. aber der Zugriff auf ein array ist überhaupt nicht vorteilhaft. In anderen Worten, ich kann nur sagen ein warp muss der Zugriff auf eine einzelne Adresse, um die positiven Optimierung/Beschleunigung, die Gewinne aus der Konstanten-Speicher zugreifen. Ist das richtig?

Außerdem möchte ich wissen, wenn ich halten eine Struktur, anstatt eine einfache Art, die in meiner ständigen Erinnerung. Jeder Zugriff auf die Struktur, die durch ein Gewinde mit einem warp; wird auch als single-memory-access-oder mehr? Ich meine eine Struktur enthalten könnte mehrere einfache Typen und array-zum Beispiel, wenn der Zugriff auf diese einfachen Typen sind, sind diese Zugriffe auch serialisiert oder nicht?

Letzte Frage wäre, falls ich ein array mit Konstanten Werten, die zugegriffen werden muss über verschiedene threads in einem warp; für schnelleren Zugriff es sollte gehalten werden im globalen Speicher anstelle von Konstanten Speicher. Ist das richtig?

Jeder kann finden Sie mir einige Beispiel-code, wo eine effiziente, Konstante Speicher Auslastung angezeigt wird.

Grüße,

InformationsquelleAutor Psypher | 2013-08-02

  1. 22

    Kann ich sagen ein warp muss der Zugriff auf eine einzelne Adresse, um die positiven Optimierung/Beschleunigung, die Gewinne aus der Konstanten-Speicher zugreifen. Ist das richtig?

    Ja, das ist in der Regel richtig und das ist die Hauptsache Absicht der Verwendung von Konstanten Speicher - /Konstanten-cache. Die Konstante cache dienen kann, bis ein 32-bit-Menge pro Zyklus pro SM. Daher, wenn jeder thread in einer Kette ist der Zugriff auf die gleichen Wert:

    int i = array[20];

    dann haben Sie die Gelegenheit, für eine gute profitieren von der ständigen cache/Speicher. Wenn jeder thread in einer Kette ist, die Zugriff auf eine einzigartige Menge:

    int i = array[threadIdx.x];

    dann die Zugriffe werden serialisiert, und die ständige Verwendung von Daten wird enttäuschend sein, performance-wise.

    Außerdem möchte ich wissen, wenn ich halten eine Struktur, anstatt eine einfache Art, die in meiner ständigen Erinnerung. Jeder Zugriff auf die Struktur, die durch ein Gewinde mit einem warp; wird auch als single-memory-access oder mehr?

    Können Sie sicher Strukturen in steter Erinnerung. Die gleichen Regeln gelten:

    int i = constant_struct_ptr->array[20];

    hat die Möglichkeit zu nutzen, aber

    int i = constant_struct_ptr->array[threadIdx.x];

    nicht. Wenn Sie auf die gleichen einfache Art Strukturelement mehrere threads, das ist ideal für Konstante cache-Nutzung.

    Letzte Frage wäre, falls ich ein array mit Konstanten Werten, die zugegriffen werden muss über verschiedene threads in einem warp; für schnelleren Zugriff es sollte gehalten werden im globalen Speicher anstelle von Konstanten Speicher. Ist das richtig?

    Ja, wenn Sie wissen, dass Sie in der Regel Ihre Zugriffe Bruch der Konstanten-Speicher, eine 32-bit-Menge pro Zyklus, Regel, dann werden Sie wahrscheinlich besser dran, verlassen die Daten in gewöhnlichen, globalen Speicher.

    Gibt es eine Vielzahl von cuda Beispiel-codes, die zeigen, dass die Verwendung von __constant__ Daten. Hier sind ein paar:

    1. Grafik volumeRender
    2. imaging bilateralFilter
    3. imaging convolutionTexture
    4. Finanzen MonteCarloGPU

    und es gibt andere.

    EDIT: Antwort auf eine Frage in die Kommentare, wenn wir eine Struktur haben, wie dies in steter Erinnerung:

    struct Simple { int a, int b, int c} s;

    Und wir wie dieser:

    int p = s.a + s.b + s.c;
          ^     ^     ^
          |     |     |
cycle:    1     2     3

    Wir haben gute Verwendung der Konstanten-Speicher/cache. Wenn der C-code kompiliert wird, unter der Haube wird es erzeugt Maschinencode, der greift auf die entsprechenden 1,2,3 in der Abbildung oben. Stellen wir uns vor, dass der Zugang 1 Eintritt. Da der Zugang 1 ist die gleichen Speicherstelle unabhängig von dem thread in den warp, während des Zyklus 1 ist, werden alle threads erhalten den Wert in s.a und er nutzt den cache optimal nutzen. Ebenso für Zugriffe 2 und 3. Wenn auf der anderen Seite wir hatten:

    struct Simple { int a[32], int b[32], int c[32]} s;
...
int idx = threadIdx.x + blockDim.x * blockIdx.x;
int p = s.a[idx] + s.b[idx] + s.c[idx];

    Dies würde nicht geben eine gute Nutzung von Konstanten-Speicher/cache. Statt, wenn diese waren typisch für unsere Zugriffe auf s hätten, hätten wir vermutlich bessere Leistung Auffinden s im gewöhnlichen globalen Speicher.

    Große und klare Antwort. Können Sie auch hinzufügen, nur ein wenig über Strukturen. z.B. meine Konstante Struktur sieht aus wie struct Simple { int a, int b, int c}. Ich bin Zugriff auf diese Typen einfach eine nach der anderen, z.B. ` p = en.a + s.b + s.c ` und dieser code wird ausgeführt, indem alle threads in einem warp; in diesem Fall werden die Zugriffe auf diese Variablen serialisiert werden oder nicht?

    InformationsquelleAutor Robert Crovella

Schreibe einen Kommentar