CUDA global (wie in C) dynamischen Arrays Gerätespeicher zugewiesen

Question

Also, im Versuch, einen Code zu schreiben, die Nvidias CUDA-Architektur nutzt. Ich bemerkte, dass nach und von dem Gerät zu kopieren war wirklich meine Gesamtleistung zu verletzen, so jetzt versuche ich, eine große Menge an Daten auf das Gerät zu bewegen.

Da diese Daten in einer Vielzahl von Funktionen verwendet werden, würde ich es gerne global sein. Ja, ich kann Zeiger laufen um, aber ich würde wirklich gerne wissen, wie mit Globals in diesem Fall zu arbeiten.

Also, ich habe Gerätefunktionen, die ein Gerät zugewiesen Array zugreifen möchten.

Idealerweise kann ich so etwas wie tun:

__device__ float* global_data;

main()
{
  cudaMalloc(global_data);
  kernel1<<<blah>>>(blah); //access global data
  kernel2<<<blah>>>(blah); //access global data again
}

Allerdings habe ich habe herausgefunden, wie ein dynamisches Array zu erstellen. Ich dachte, eine Arbeit um, indem das Array deklariert wie folgt:

__device__ float global_data[REALLY_LARGE_NUMBER];

Und während das kein cudaMalloc Anruf erforderlich ist, würde ich die dynamische Zuweisung Ansatz bevorzugen.

Answer 1

So etwas sollte wahrscheinlich funktionieren.

#include <algorithm>

#define NDEBUG
#define CUT_CHECK_ERROR(errorMessage) do {                                 \
        cudaThreadSynchronize();                                           \
         cudaError_t err = cudaGetLastError();                             \
         if( cudaSuccess != err) {                                         \
                     fprintf(stderr, "Cuda error: %s in file '%s' in line %i : %s.\n",    \
                                             errorMessage, __FILE__, __LINE__, cudaGetErrorString( err) );\
                     exit(EXIT_FAILURE);                                                  \
                 } } while (0)


__device__ float *devPtr;

__global__
void kernel1(float *some_neat_data)
{
    devPtr = some_neat_data;
}

__global__
void kernel2(void)
{
    devPtr[threadIdx.x] *= .3f;
}


int main(int argc, char *argv[])
{
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    CUT_CHECK_ERROR("kernel1");

    kernel2<<<1,128>>>();

    CUT_CHECK_ERROR("kernel2");

    return 0;
}

Geben Sie ihm einen Wirbel.

Answer 2

Verbringen Sie einige Zeit auf der reichlichen Dokumentation konzentriert durch NVIDIA angeboten.

Aus dem Programmierhandbuch:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

Das ist ein einfaches Beispiel, wie Speicher zuzuweisen. Nun, in Ihrem Kernel, sollten Sie einen Zeiger auf einen Schwimmer wie so akzeptieren:

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x]++;
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= 0.3f;
}

So, jetzt

Sie können sie aufrufen, etwa so:

float* devPtr;
cudaMalloc((void**)&devPtr, 256 * sizeof(*devPtr));
cudaMemset(devPtr, 0, 256 * sizeof(*devPtr));

kernel1<<<1,128>>>(devPtr);
kernel2<<<1,128>>>(devPtr);

  

Da diese Daten in zahlreichen verwendet wird   Funktionen, würde ich es gerne sein   global.

Es gibt einige gute Gründe, Globals zu verwenden. Dies ist definitiv nicht ein. Ich werde es als eine Übung verlassen, um dieses Beispiel zu erweitern um „devPtr“ zu einem globalen Rahmen zu bewegen.

EDIT:

Ok, das grundlegende Problem ist dies: Ihr Kernel kann nur Speicherzugriffsgerät und die einzigen globalen Umfang Hinweise, die sie sind GPU diejenigen nutzen können. Wenn ein Kernel aus Ihrer CPU ruft, hinter den Kulissen, was passiert ist, dass die Zeiger und Primitiven in GPU-Register kopiert werden und / oder die gemeinsamen Speicher, bevor der Kernel ausgeführt wird.

So engster kann ich vorschlagen, ist dies: verwenden cudaMemcpyToSymbol (), um Ihre Ziele zu erreichen. Aber im Hintergrund, die Ansicht, dass ein anderer Ansatz könnte das Richtige sein.

#include <algorithm>

__constant__ float devPtr[1024];

__global__
void kernel1(float *some_neat_data)
{
    some_neat_data[threadIdx.x] = devPtr[0] * devPtr[1];
}

__global__
void kernel2(float *potentially_that_same_neat_data)
{
    potentially_that_same_neat_data[threadIdx.x] *= devPtr[2];
}


int main(int argc, char *argv[])
{
    float some_data[256];
    for (int i = 0; i < sizeof(some_data) / sizeof(some_data[0]); i++)
    {
        some_data[i] = i * 2;
    }
    cudaMemcpyToSymbol(devPtr, some_data, std::min(sizeof(some_data), sizeof(devPtr) ));
    float* otherDevPtr;
    cudaMalloc((void**)&otherDevPtr, 256 * sizeof(*otherDevPtr));
    cudaMemset(otherDevPtr, 0, 256 * sizeof(*otherDevPtr));

    kernel1<<<1,128>>>(otherDevPtr);
    kernel2<<<1,128>>>(otherDevPtr);

    return 0;
}

Nicht vergessen '--host-Kompilation = c ++' für dieses Beispiel.

Answer 3

Ich ging weiter und versuchte, die Lösung eines temporären Zeiger Aufteilung und es auf eine einfache globale Funktion ähnlich kernel1 vorbei.

Die gute Nachricht ist, dass es funktioniert:)

Aber ich denke, es den Compiler verwirrt wie ich jetzt „Advisory: kann nicht sagen, welche Zeiger zeigt auf, globalen Speicherplatz unter der Annahme,“ bekommen, wenn ich versuche, die globalen Daten zuzugreifen. Glücklicherweise geschieht die Annahme richtig zu sein, aber die Warnungen sind ärgerlich.

Wie auch immer, für den Rekord - ich habe schon in vielen der Beispiele gesucht und haben durch die nvidia Übungen laufen, wo der Punkt, um die Ausgabe zu erhalten, ist zu sagen: „Richtig!“. Allerdings habe ich nicht in alle von ihnen sieht. Wenn jemand eines sdk Beispiel weiß, wo sie dynamische globale Gerätespeicherzuweisung zu tun, würde ich noch gerne wissen.

Answer 4

Erm, es war genau das Problem der devPtr zu globalem Bereich zu bewegen, das mein Problem war.

Ich habe eine Implementierung, die genau das tut, mit den beiden Kernen in einen Zeiger auf Daten übergeben zu müssen. Ich ausdrücklich nicht will, in dem Zeiger zu übergeben.

ich die Dokumentation ziemlich genau gelesen habe, und die nvidia-Foren treffen (und Google für eine Stunde gesucht oder so), aber ich habe nicht eine Implementierung einer globalen und dynamischen Vorrichtungsanordnung, die tatsächlich läuft (ich habe versucht, mehr gefunden dass kompiliert und dann in neuen und interessanten Möglichkeiten gescheitert).

Answer 5

überprüfen Sie die Proben mit dem SDK enthalten. Viele dieser Beispielprojekte sind eine anständige Art und Weise durch das Beispiel zu lernen.

Answer 6

  

Da diese Daten in einer Vielzahl von Funktionen verwendet wird, würde ich es gerne global sein.

-

  

Es gibt einige gute Gründe, Globals zu verwenden. Dies ist definitiv nicht ein. Ich werde es wie ein verlassen   Übung dieses Beispiel zu erweitern „devPtr“ zu einem globalen Rahmen schließen bewegen.

Was passiert, wenn der Kernel auf einem großen konst Struktur arbeitet von Arrays aus? die so genannten Konstante Speicher verwendet, ist keine Option, weil es in der Größe sehr begrenzt ist .. so ist, dann müssen Sie es im globalen Speicher setzen ..?