CUDA设备端复杂对象初始化的cudaDeviceSetLimit正确配置问题

• 你没有计算设备端new操作的元数据开销：CUDA设备堆上的每次动态分配都会额外占用16~64字节不等的管理内存，用于存储分配块长度、对齐信息等元数据，你仅计算了对象本身的内存占用，没有包含这部分开销。
• cudaDeviceSetLimit调用顺序错误：cudaLimitMallocHeapSize必须在程序发起任何CUDA API调用（包括cudaMalloc、上下文初始化相关操作）之前执行，你在cudaMalloc之后才设置堆大小，配置完全不生效，仅靠默认的8MB设备堆容量支撑，块数少的时候刚好够，块数多了直接溢出。
• 你的主机端代码存在无效内存申请：先在主机堆上new了arr_gpu_net指针数组，马上用cudaMalloc覆盖了该指针，造成主机内存泄漏。

1. 调整堆大小计算逻辑，在你计算的基础值上预留至少30%的冗余空间，覆盖元数据和对齐开销
2. 把cudaDeviceSetLimit调用移到所有CUDA操作之前
3. 如果后续要给类新增成员，同步更新基础大小的计算逻辑，可适当提高冗余系数

class NNode{
public:
    int node_id;
};

class cuNetwork{
public:
    int num_allnodes;
    NNode** all_nodes; 
};

__global__ void mallocTest(int num_allnodes, cuNetwork** arr_gpu_net){
    int bId = blockIdx.x; 
    cuNetwork* gpu_net  = new cuNetwork(); 
    gpu_net->all_nodes = new NNode*[num_allnodes];
    for(int i=0; i<num_allnodes; i++){
        gpu_net->all_nodes[i] = new NNode();
    }
    arr_gpu_net[bId] = gpu_net;
}

int main(int argc, const char **argv){
    int numBlocks = 65535; 
    int num_allnodes = 200; 

    // 最先设置设备堆大小
    size_t base_size = sizeof(cuNetwork) + sizeof(NNode*) * num_allnodes + sizeof(NNode) * num_allnodes;
    // 预留50%冗余覆盖元数据、对齐开销，复杂对象可提高到2倍
    size_t total_heap_size = (size_t)(numBlocks * base_size * 1.5);
    // 最小不低于CUDA默认的8MB，避免小负载下出现异常
    if(total_heap_size < 8 * 1024 * 1024) total_heap_size = 8 * 1024 * 1024;
    cudaDeviceSetLimit(cudaLimitMallocHeapSize, total_heap_size);

    // 再执行CUDA内存分配等其他操作
    cuNetwork** arr_gpu_net;
    cudaMalloc((void **)&arr_gpu_net, sizeof(cuNetwork*) * numBlocks);

    mallocTest<<<numBlocks, 1>>>(num_allnodes, arr_gpu_net);
    cudaDeviceSynchronize();

    // 业务逻辑执行完成后记得释放设备堆内存，避免显存泄漏
    return 0;
}

额外提醒：如果调整后仍报错，可先检查cudaDeviceSetLimit的返回值确认是否设置成功，同时确认当前设备的剩余显存大于你设置的堆大小。

内容的提问来源于stack exchange，提问作者Joachim