向CUDA Kernel高效传递参数时如何减少cudaMalloc和cudaMemcpy调用

以下是几种可大幅减少冗余CUDA API调用的实现方式，均满足传递结构体指针而非拷贝的要求：

方案1：使用CUDA统一内存（最高效精简代码） #

CUDA的统一内存（Unified Memory）机制提供主机和设备共享的地址空间，无需手动维护主机、设备两份数据副本，也不需要手动执行内存拷贝，代码量可缩减一半以上，所有开普勒及以上架构的CUDA设备均支持该特性。
示例实现：

struct Data {
    float* a;
    float* b;
    float* c;
    // 其余指针成员
};

__global__ void foo(Data* data, size_t n) {
    size_t const id = blockIdx.x * blockDim.x + threadIdx.x;
    if (id < n) {
        data->c[id] = data->a[id] + data->b[id];
    }
}

int main() {
    const size_t kSize = 10;
    // 统一内存分配Data结构体，主机设备均可直接访问
    Data* data;
    cudaMallocManaged(&data, sizeof(Data));
    // 数组也分配在统一内存空间
    cudaMallocManaged(&data->a, sizeof(float) * kSize);
    cudaMallocManaged(&data->b, sizeof(float) * kSize);
    cudaMallocManaged(&data->c, sizeof(float) * kSize);

    // 主机端直接初始化数据，无需额外拷贝
    for (size_t i = 0; i < kSize; ++i) {
        data->a[i] = 10.f;
        data->b[i] = 10.f;
    }

    // 直接调用内核，无需手动同步结构体和数组数据
    const size_t block_size = 512;
    const size_t grid_size = (kSize + block_size - 1) / block_size;
    foo<<<grid_size, block_size>>>(data, kSize);

    cudaDeviceSynchronize();

    // 主机端直接读取计算结果
    for (size_t i = 0; i < kSize; ++i) {
        printf("%f\n", data->c[i]);
    }

    // 统一释放内存
    cudaFree(data->a);
    cudaFree(data->b);
    cudaFree(data->c);
    cudaFree(data);
    return 0;
}

方案2：打包连续内存块减少API调用次数 #

如果不使用统一内存，可以将所有数组合并为一块连续的设备内存，同时将结构体的批量成员拷贝改为单次全量拷贝，大幅降低API调用开销：

• 所有数组仅需1次cudaMalloc分配连续内存，通过偏移量区分不同数组成员
• 整个结构体仅需1次cudaMemcpy即可完成所有设备指针的绑定，替代原先逐成员拷贝的操作
• 如果主机端也将数组存储为连续内存块，所有数据的主机到设备拷贝也可合并为1次调用

示例核心逻辑：

const size_t kSize = 10;
const size_t total_arr_bytes = 3 * kSize * sizeof(float);
// 一次分配所有数组的连续设备内存
float* dev_buffer;
cudaMalloc(&dev_buffer, total_arr_bytes);

// 给设备端结构体的数组成员赋值偏移地址
struct Data dev_data;
dev_data.a = dev_buffer;
dev_data.b = dev_buffer + kSize;
dev_data.c = dev_buffer + 2 * kSize;

// 一次拷贝整个结构体到设备，替代逐成员拷贝
struct Data* dev_data_ptr;
cudaMalloc(&dev_data_ptr, sizeof(Data));
cudaMemcpy(dev_data_ptr, &dev_data, sizeof(Data), cudaMemcpyHostToDevice);

// 主机端如果用连续内存存储数组，可一次拷贝所有数据到设备
float host_buffer[3 * kSize];
for (size_t i = 0; i < kSize; ++i) {
    host_buffer[i] = 10.f; // a数组
    host_buffer[i + kSize] = 10.f; // b数组
}
cudaMemcpy(dev_buffer, host_buffer, total_arr_bytes, cudaMemcpyHostToDevice);

// 内核调用逻辑不变
foo<<<grid_size, block_size>>>(dev_data_ptr, kSize);

方案3：封装辅助宏/工具函数 #

不管采用上述哪种方案，都可以通过封装重复逻辑进一步减少冗余代码，比如定义宏批量处理成员的分配、拷贝、释放操作：

// 批量分配设备成员内存的宏
#define ALLOC_DEV_MEMBER(dev_struct, member, bytes) \
    cudaMalloc((void**)&dev_struct.member, bytes)
// 批量拷贝数据到设备的宏
#define COPY_TO_DEV(dev_struct, host_struct, member, bytes) \
    cudaMemcpy(dev_struct.member, host_struct.member, bytes, cudaMemcpyHostToDevice)
// 同理可定义拷贝回主机、释放内存的对应宏

避免逐成员手写重复的API调用，同时降低出错概率。

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