如何计算内核中动态内存分配所需的堆大小？

Question

我面临的问题是，如果我将 CUDA 堆大小设置为我需要在内核中分配的总内存量，那么堆仍然不足以满足所有分配。

这是一个代表我的用例的最小示例：

#include <stdio.h>

#define NARR 8

__global__
void heaptest(int N){
    double* arr[NARR];

    __shared__ double* arrS[NARR];

    if(threadIdx.x == 0){
        for(int i = 0; i < NARR; i++){
            arrS[i] = (double*) malloc(sizeof(double) * N);

            if(arrS[i] == NULL)
                printf("block %d, array %d is NULL\n", blockIdx.x, i);
        }            
    }

    __syncthreads();

    for(int i = 0; i < NARR; i++){
        arr[i] = arrS[i];
    }
}

size_t getHeapSizePerBlock(int N){
    return sizeof(double) * N * NARR;
}

int main(){

    int N = 4000 * 18;

    int nBlocks = 1;

    size_t myheapsize = getHeapSizePerBlock(N) * nBlocks;

    printf("set heap size to %lu\n", myheapsize);
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, myheapsize);

    size_t a;

    cudaDeviceGetLimit(&a, cudaLimitMallocHeapSize);
    printf("heap size is now %lu\n", a);

    heaptest<<<nBlocks, 128>>>(N);

    cudaDeviceSynchronize();

    cudaDeviceReset();

    return 0;
}

我用 nvcc V8.0.61 编译。

nvcc -arch=sm_60 heaptest.cu -o heaptest

程序输出是

set heap size to 4608000
heap size is now 4653056
block 0, array 7 is NULL

因此，即使堆大小大于所需大小，也不够大。 在这种情况下如何正确计算所需的尺寸？

Answer 1

您可能无法计算出应用程序堆所需的准确大小，因为您无法控制 CUDA 的内存管理器。就像在拥有操作系统内存管理器的地方分配 CPU 内存一样，CUDA 也有自己的内存管理器。当您在堆中分配多个数组时，您无法保证它们完全适合堆的大小，可能存在一些开销。

举例来说，我对您的代码做了一个小的修改，以打印 malloc 返回的内存地址：

printf("block %d, array %d is %p\n", blockIdx.x, i, arrS[i]);

这是我在 GTX 1070 上得到的：

block 0, array 0 is 0x102059a8d20
block 0, array 1 is 0x10205600120
block 0, array 2 is 0x1020568f280
block 0, array 3 is 0x10205738520
block 0, array 4 is 0x102057c7680
block 0, array 5 is 0x10205870920
block 0, array 6 is 0x102058ffa80
block 0, array 7 is (nil)

首先要注意的是内存位置不是（总是）连续/增加的（例如，数组 0 > 数组 6 > ... > 数组 1），但这对我们来说并不重要。此外，如果按降序减去内存地址，则不会得到传递给malloc() 的大小，在您的情况下始终为sizeof(double) * N 或576000 字节。例如：

0x1020568f280 - 0x10205600120 = 586080 字节（数组 1）

0x10205738520 - 0x1020568f280 = 692896 字节（数组 2）

由于相对于传递给malloc() 的块大小而言，这些块在内存中是连续的，因此我们可以验证确实存在一些我们无法分配576000 字节块的内存块。在数组 1 和 2 之间，我们有额外的 10080 字节，在数组 2 和 3 之间，有额外的 116896 字节（这超过了块大小的 20%！）。

我要做的是避免在堆上动态分配内存，而是在主机代码执行期间分配它。但是，如果您出于某种原因确实需要这样做，我建议通过在似乎足够之前对其进行测试来设置具有一些开销余量的堆大小。我至少希望即使存在一些堆分配开销，这也不应该太大，所以如果需要，可以开始分配额外的 10% 并从那里增加。