【CUDA】カーネル内での動的メモリ確保

バイス上のメモリは、ホストからのcudaMalloc()で確保するのが通常だが、 カーネル内で動的にメモリ確保することもできる*1カーネル内での動的メモリ確保には、C/C++と同じようにmalloc()/free()new/deleteが使える。

__global__
void testKernel() {
  size_t size = 1024;

  // Dynamic allocation with malloc/free
  char *a = (char*)malloc(size);
  free(a);

  // Dynamic allocation with new/delete
  char *b = new char[size];
  delete[] b;
}

int main() {
  testKernel<<<1, 1>>>();
  cudaDeviceSynchronize();
  return 0;
}

ただし、ホストからのcudaMalloc()カーネル内でのmalloc()はメモリ確保方法が異なる。 ホストからのcudaMalloc()はデバイスメモリの未使用の領域にメモリを確保するのに対して、 カーネル内でのmalloc()はあらかじめ用意されたデバイス上のヒープにメモリを確保する。 そのため、ホストからのcudaMalloc()で確保したメモリはカーネル内では解放できないし、 カーネル内で確保したメモリはホストからは解放できない。 以下のコードはエラーになるパターンを表している。

__global__
void mallocInKernel(char **ptr, size_t size) {
  *ptr = (char*)malloc(size);
}

__global__
void freeInKernel(char *ptr) {
  free(ptr);
}

int main() {
  size_t size = 1024;

  // ホスト上でcudaMalloc()したメモリをカーネル内でfree()
  char *a;
  cudaMalloc(&a, size);
  freeInKernel<<<1, 1>>>(a);  // Error!!
  cudaDeviceSynchronize();

  // カーネル内でmalloc()したメモリをホストからcudaFree()
  char *hb, **db;
  cudaMalloc(&db, sizeof(char*));
  mallocInKernel<<<1, 1>>>(db, size);
  cudaMemcpy(&hb, db, sizeof(char*), cudaMemcpyDeviceToHost);
  cudaFree(hb);  // Error!!
  cudaDeviceSynchronize();

  return 0;
}

バイスのヒープはデフォルトで8MBしか確保されない。 つまり、カーネル内のmalloc()で8MB以上のメモリを確保するとエラーになる。 この設定を変えるにはcudaDeviceSetLimit(cudaLimitMallocHeapSize, size_t size)を使えばよい。 例えば、以下のように書くと1GBまでカーネル内でmalloc()できる。

size_t size = 1024 * 1024 * 1024;
cudaDeviceSetLimit(cudaLimitMallocHeapSize, size);

この設定は全てのカーネル実行が行われる前に書かれている必要があり、途中で設定を変えることはできない。 また、現在のヒープサイズの設定がどうなっているか確認するためには cudaDeviceGetLimit(size_t* size, cudaLimitMallocHeapSize)を使えばよい。

size_t size;
cudaDeviceGetLimit(&size, cudaLimitMallocHeapSize);
printf("Heap Size=%ld\n", size);

Compute Capability 3.5以降ではカーネル内でcudaMalloc()/cudaFree()を使うことができるが、 これはカーネル内でのmalloc()/free()と同じ効果をもたらす

バイスのヒープサイズを無駄に拡げないためにも、できるだけホストからのcudaMalloc()で メモリを確保したいが、場合によってはカーネル内での動的メモリ確保が有効なケースもあるだろう。

*1:Compute Capability 2.x以降でのみサポート