【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以降でのみサポート