20

カーネル関数内でいくつかの配列を動的に割り当てる必要があります。どうすればそれができますか？

私のコードは次のようなものです：

__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 分間実行されるためです。

4

5 に答える 5

14

@talonmies が、カーネル内でメモリを動的に割り当てる方法に関する質問に回答しました。これは、補足的な回答として意図されており、のパフォーマンス__device__ malloc()と、検討したい代替案に対処しています。

カーネルでメモリを動的に割り当てると、GPU コードが CPU コードのように見えるため、魅力的な場合があります。しかし、パフォーマンスに深刻な影響を与える可能性があります。自己完結型のテストを作成し、以下に含めました。このテストでは、約 260 万のスレッドが起動されます。各スレッドは、グローバル メモリの 16 個の整数にスレッド インデックスから派生した値を設定し、値を合計して合計を返します。

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

私の 2.0 デバイスでは、カーネルは事前割り当てメモリを使用する__device__ malloc()と 1500 ミリ秒で実行され、事前割り当てメモリを使用すると 27 ミリ秒で実行されます。つまり、メモリがカーネル内で動的に割り当てられると、テストの実行に 56 倍の時間がかかります。時間には、カーネルの一部ではない外側のループ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;
}

出力：

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
于 2012-11-21T02:08:41.427 に答える