如何在內核中動態分配數組？

Question

我需要在內核函數內動態分配一些數組。我怎麼能這樣做？

我的代碼是類似的東西：

__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. 
} 

但是，這是行不通的。如果這是在主機代碼中，我可以使用malloc。 cudaMalloc需要主機上的指針，以及其他設備上的指針。在內核函數中，我沒有主機指針。

那麼，我該怎麼辦？

如果分配所有數組需要很長時間（我需要大約4個大小n和5個大小nn），這不會是一個問題。由於內核可能會運行至少20分鐘。

Answer 1

只有計算能力2.x和更新的硬件才支持動態內存分配。您可以使用C++的新關鍵字或malloc的內核，所以你的例子可以變成：

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

這在具有上下文的生命週期本地內存運行時堆分配內存，所以一定要釋放如果你的意圖不是再次使用內存，內核完成運行後的內存。您還應該注意，無法直接從主機API訪問運行時堆內存，因此，例如，您無法將分配給內核的指針作爲參數傳遞給cudaMemcpy。

Answer 2

@talonmies回答了你如何在內核中動態分配內存的問題。這旨在作爲補充答案，解決__device__ malloc()的性能問題以及您可能需要考慮的替代方案。

在內核中動態分配內存很誘人，因爲它允許GPU代碼更像CPU代碼。但它會嚴重影響性能。我寫了一個自包含的測試，並將其包含在下面。測試啓動大約260萬個線程。每個線程使用從線索索引派生的一些值填充全局內存的16個整數，然後總結這些值並返回總和。

該測試實現了兩種方法。第一種方法使用__device__ malloc()，第二種方法使用內核運行之前分配的內存。

在我的2.0設備上，當使用__device__ malloc()時，內核在1500ms內運行，使用預先分配的內存時內核運行時間爲27ms。換句話說，在內核中動態分配內存時，測試需要56x更長的時間運行。時間包括外部循環cudaMalloc()/cudaFree()，它不是內核的一部分。如果多次使用相同數量的線程啓動相同的內核（通常情況如此），則在所有內核啓動時分攤成本。這個差距甚至更高，達到60x左右。

推測，我認爲性能衝擊部分是由隱式序列化引起的。 GPU可能必須序列化所有同時調用__device__ malloc()，以便爲每個調用者提供單獨的內存塊。

在運行內核之前，不使用__device__ malloc()的版本將分配所有GPU內存。內存指針傳遞給內核。每個線程計算一個索引到先前分配的內存中，而不是使用__device__ malloc()。

預先分配內存的潛在問題是，如果只有某些線程需要分配內存，並且不知道哪些線程是哪些線程，則有必要爲所有線程分配內存。如果沒有足夠的內存，那麼使用__device__ malloc()減少每個內核調用的線程數可能更有效。其他解決方法可能最終會重新實現__device__ malloc()在後臺執行的操作，並且會看到類似的性能下降。

測試__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; 
} 

輸出：

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

如果n的值和NN被稱爲內核調用之前，那麼爲什麼不cudaMalloc在主機端內存並將設備內存指針傳遞給內核？