Wie man Arrays innerhalb eines Kernels dynamisch zuordnet?

Question

Ich muss einige Arrays innerhalb der Kernel-Funktion dynamisch zuweisen. Wie kann ich das tun?

Mein Code ist so etwas wie das:

__global__ func(float *grid_d,int n, int nn){ 
    int i,j; 
    float x[n],y[nn]; 
    //Do some really cool and heavy computations here that takes hours. 
} 

Aber das wird nicht funktionieren. Wenn dies innerhalb des Host-Codes war, könnte ich malloc verwenden. cudaMalloc benötigt einen Zeiger auf den Host und andere auf dem Gerät. Innerhalb der Kernel-Funktion habe ich den Host-Zeiger nicht.

Also, was soll ich tun?

Wenn es zu lange dauert (einige Sekunden), um alle Arrays zuzuordnen (ich brauche etwa 4 der Größe n und 5 der Größe nn), wird dies kein Problem sein. Da der Kernel wird wahrscheinlich für mindestens 20 Minuten laufen.

Answer 1

Die dynamische Speicherzuweisung wird nur auf der Compute-Fähigkeit 2.x und neuer Hardware unterstützt. Sie können entweder das C++ neues Schlüsselwort oder malloc im Kernel verwenden, so dass Ihr Beispiel werden könnte:

__global__ func(float *grid_d,int n, int nn){ 
    int i,j; 
    float *x = new float[n], *y = new float[nn]; 
} 

Dieser Speicher auf einem lokalen Speicher Runtime Heap reserviert, die die Lebensdauer des Kontextes hat, so stellen Sie sicher, dass Sie die kostenlose Speicher, nachdem der Kernel ausgeführt wurde, wenn Sie den Speicher nicht erneut verwenden möchten. Sie sollten außerdem beachten, dass auf den Heapspeicher im Runtime nicht direkt von den Host-APIs zugegriffen werden kann. Daher können Sie beispielsweise keinen Zeiger übergeben, der in einem Kernel als Argument an cudaMemcpy zugewiesen ist.

Answer 2

@talonmies beantwortet Ihre Frage zur dynamischen Speicherzuweisung in einem Kernel. Dies ist als ergänzende Antwort gedacht, die sich mit der Leistung von __device__ malloc() und einer Alternative, die Sie in Betracht ziehen könnten, befasst.

Das dynamische Zuordnen von Speicher im Kernel kann verlockend sein, da GPU-Code mehr wie CPU-Code aussieht. Aber es kann die Leistung ernsthaft beeinträchtigen. Ich habe einen eigenständigen Test geschrieben und ihn unten eingeschlossen. Der Test startet etwa 2,6 Millionen Threads. Jeder Thread füllt 16 Ganzzahlen des globalen Speichers mit einigen Werten, die aus dem Thread-Index abgeleitet werden, addiert dann die Werte und gibt die Summe zurück.

Der Test implementiert zwei Ansätze. Der erste Ansatz verwendet __device__ malloc() und der zweite Ansatz verwendet Speicher, der zugewiesen wird, bevor der Kernel ausgeführt wird.

Auf meinem Gerät 2.0 läuft der Kernel in 1500ms bei Verwendung __device__ malloc() und 27ms bei der Verwendung von vor-zugewiesenen Speicher. Mit anderen Worten, der Test dauert 56x länger, um ausgeführt zu werden, wenn Speicher dynamisch innerhalb des Kernels zugewiesen wird. Die Zeit enthält die äußere Schleife cudaMalloc()/cudaFree(), die nicht Teil des Kernels ist. Wenn der gleiche Kernel mehrmals mit der gleichen Anzahl von Threads gestartet wird, wie es oft der Fall ist, amortisieren sich die Kosten des cudaMalloc()/cudaFree() über alle Kernel-Starts. Das bringt den Unterschied sogar auf etwa 60x.

Speculating, ich denke, dass der Leistungseinbruch teilweise durch implizite Serialisierung verursacht wird. Die GPU muss wahrscheinlich alle gleichzeitigen Aufrufe an __device__ malloc() serialisieren, um jedem Anrufer separate Speicherblöcke zur Verfügung zu stellen.

Die Version, die __device__ malloc() nicht verwendet, weist den gesamten GPU-Speicher zu, bevor der Kernel ausgeführt wird. Ein Zeiger auf den Speicher wird an den Kernel übergeben. Jeder Thread berechnet einen Index in den zuvor zugewiesenen Speicher anstelle von __device__ malloc().

Das potenzielle Problem bei der Speicherzuordnung ist, dass, wenn nur einige Threads Speicher reservieren müssen und es nicht bekannt ist, welche Threads diese sind, es notwendig sein wird, Speicher für alle Threads zuzuweisen. Wenn nicht genügend Arbeitsspeicher dafür vorhanden ist, kann es effizienter sein, die Anzahl der Threads pro Kernelaufruf zu reduzieren, indem Sie __device__ malloc() verwenden. Andere Problemumgehungen würden wahrscheinlich dazu führen, dass das, was __device__ malloc() tut, im Hintergrund reimplementiert wird und einen ähnlichen Leistungseinbruch sehen würde.

Testen Sie die Leistung von __device__ malloc():

#include "cuda_runtime.h" 
#include "device_launch_parameters.h" 
#include <stdio.h> 

const int N_ITEMS(16); 

#define USE_DYNAMIC_MALLOC 

__global__ void test_malloc(int* totals) 
{ 
    int tx(blockIdx.x * blockDim.x + threadIdx.x); 

    int* s(new int[N_ITEMS]); 

    for (int i(0); i < N_ITEMS; ++i) { 
    s[i] = tx * i; 
    } 

    int total(0); 
    for (int i(0); i < N_ITEMS; ++i) { 
    total += s[i]; 
    } 

    totals[tx] = total; 

    delete[] s; 
} 

__global__ void test_malloc_2(int* items, int* totals) 
{ 
    int tx(blockIdx.x * blockDim.x + threadIdx.x); 

    int* s(items + tx * N_ITEMS); 

    for (int i(0); i < N_ITEMS; ++i) { 
    s[i] = tx * i; 
    } 

    int total(0); 
    for (int i(0); i < N_ITEMS; ++i) { 
    total += s[i]; 
    } 

    totals[tx] = total; 
} 

int main() 
{ 
    cudaError_t cuda_status; 

    cudaSetDevice(0); 

    int blocks_per_launch(1024 * 10); 
    int threads_per_block(256); 

    int threads_per_launch(blocks_per_launch * threads_per_block); 

    int* totals_d; 
    cudaMalloc((void**)&totals_d, threads_per_launch * sizeof(int)); 

    cudaEvent_t start, stop; 
    cudaEventCreate(&start); 
    cudaEventCreate(&stop); 

    cudaDeviceSynchronize(); 
    cudaEventRecord(start, 0); 

#ifdef USE_DYNAMIC_MALLOC 
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, threads_per_launch * N_ITEMS * sizeof(int)); 

    test_malloc<<<blocks_per_launch, threads_per_block>>>(totals_d); 
#else 
    int* items_d; 
    cudaMalloc((void**)&items_d, threads_per_launch * sizeof(int) * N_ITEMS); 

    test_malloc_2<<<blocks_per_launch, threads_per_block>>>(items_d, totals_d); 

    cudaFree(items_d); 
#endif 

    cuda_status = cudaDeviceSynchronize(); 
    if (cuda_status != cudaSuccess) { 
    printf("Error: %d\n", cuda_status); 
    exit(1); 
    } 

    cudaEventRecord(stop, 0); 
    cudaEventSynchronize(stop); 
    float elapsedTime; 
    cudaEventElapsedTime(&elapsedTime, start, stop); 

    printf("Elapsed: %f\n", elapsedTime); 

    int* totals_h(new int[threads_per_launch]); 
    cuda_status = cudaMemcpy(totals_h, totals_d, threads_per_launch * sizeof(int), cudaMemcpyDeviceToHost); 
    if (cuda_status != cudaSuccess) { 
    printf("Error: %d\n", cuda_status); 
    exit(1); 
    } 

    for (int i(0); i < 10; ++i) { 
    printf("%d ", totals_h[i]); 
    } 
    printf("\n"); 

    cudaFree(totals_d); 
    delete[] totals_h; 

    return cuda_status; 
} 

Ausgang:

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe 
Elapsed: 27.311169 
0 120 240 360 480 600 720 840 960 1080 

C:\rd\projects\test_cuda_malloc\Release>test_cuda_malloc.exe 
Elapsed: 1516.711914 
0 120 240 360 480 600 720 840 960 1080 

Answer 3

Wenn der Wert von n und nn bekannt war, bevor der Kern genannt wird, warum dann nicht den Speicher auf Host-Seite cudaMalloc und übergeben Sie den Gerätespeicherzeiger an den Kernel?