カーネル内で動的に配列を割り当てる方法は？

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は、ホスト上のポインタとデバイス上のポインタを必要とします。カーネル関数の中で私はホストポインタを持っていません。

どうすればよいですか？

すべての配列を割り当てるには時間がかかりすぎます（サイズnが約4、サイズがnnが約5です）。これは問題ありません。カーネルはおそらく少なくとも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個の整数にスレッドインデックスから派生したいくつかの値を設定し、その値を合計して合計を返します。

このテストでは、2つのアプローチが実装されています。最初のアプローチは__device__ malloc()を使用し、2番目のアプローチはカーネルが実行される前に割り当てられたメモリを使用します。

私の2.0デバイスでは、カーネルは__device__ malloc()を使用すると1500msで実行され、事前割り当てメモリを使用する場合は27msです。言い換えれば、メモリがカーネル内で動的に割り当てられるときに、が56x長く、が実行されます。この時間には、外部ループcudaMalloc()/cudaFree()が含まれていますが、これはカーネルの一部ではありません。同じカーネルが同じスレッド数で何度も起動された場合、よくあることですが、cudaMalloc()/cudaFree()のコストはすべてのカーネル起動で償却されます。その差はさらに大きくなり、約60倍になります。

推測すると、パフォーマンスの低下は暗黙的なシリアライゼーションによって一部引き起こされていると思います。おそらく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; 
} 

出力：nの値とNNが知られていた場合は

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

カーネルが呼び出される前に、なぜホスト側のメモリをcudaMallocありませんデバイスのメモリポインタをカーネルに渡しますか？